diff --git a/README.md b/README.md index 24b4a743..5ca13dc8 100644 --- a/README.md +++ b/README.md @@ -374,7 +374,7 @@ The standalone EP checks live under `tests/ep/` and can be built without externa ```bash source scripts/common_env.sh -bash tests/ep/build.sh source-only +bash tests/ep/build.sh host-only ctest --test-dir tests/ep/build --output-on-failure ``` diff --git a/docs/CI.md b/docs/CI.md index 1914222d..8c49e631 100644 --- a/docs/CI.md +++ b/docs/CI.md @@ -49,7 +49,7 @@ foreign process collision. | Home | `/home/tilexr-ci` | | Sealed CANN 9.1 | `/home/tilexr-ci/toolchains/cann/9.1.0` | | Sealed Bisheng compiler | `/home/tilexr-ci/toolchains/cann/9.1.0/cann/tools/bisheng_compiler/bin/bisheng` | -| Sealed controller | `/home/tilexr-ci/control/v4` | +| Sealed controller | `/home/tilexr-ci/control/v5` | | Active controller link | `/home/tilexr-ci/control/current` | | Actions runner | `/home/tilexr-ci/actions-runner` | | Runner workspace root | `/home/tilexr-ci/actions-runner/_work` | @@ -92,15 +92,14 @@ cleanup. ## Test matrix -[`host_checks.sh`](../scripts/ci/host_checks.sh) runs these eight recorded cases: +[`host_checks.sh`](../scripts/ci/host_checks.sh) runs these seven recorded cases: - shell syntax for tracked CI and affected test scripts; - the complete standalone `tests/ci` CTest suite; -- comm logging, spdlog compile, and source-guard binaries; -- all four EP source-only tests; -- both data-as-flag tests; +- comm logging and spdlog compile tests; +- EP layout and host validation tests; +- the data-as-flag header compile test; - collectives vLLM patch tests; -- collectives vLLM integration-source tests; - collectives profile-report tests. On `blue`, [`build_blue.sh`](../scripts/ci/control/build_blue.sh) performs a @@ -113,14 +112,11 @@ failure. The build phase logs these named cases: - top level: `top-level-ctest`; -- comm: `comm-log`, `comm-spdlog-compile`, and `comm-source-guards`; -- UDMA: `udma-transport-layout`, `udma-registry`, `udma-demo-sources`, and - `udma-source-guard`; +- comm: `comm-log` and `comm-spdlog-compile`; +- UDMA: `udma-transport-layout` and `udma-registry`; - SDMA: `sdma-metadata`, `sdma-api-invalid`, `sdma-transport-disabled`, - `sdma-comm-wiring`, `sdma-source-guard`, and `sdma-header-compile`; -- EP: `ep-layout`, `ep-api-sources`, `ep-kernel-sources`, and - `ep-host-validation`; -- memory: `memory-demo-sources`. + and `sdma-header-compile`; +- EP: `ep-layout` and `ep-host-validation`. The top-level configure/build, suite builds, top-level reinstall, required-file checks, and dependency/RPATH checks are also mandatory. They fail the build diff --git a/scripts/ci/control/VERSION b/scripts/ci/control/VERSION index c694117f..47e5d40a 100644 --- a/scripts/ci/control/VERSION +++ b/scripts/ci/control/VERSION @@ -1 +1 @@ -v4 +v5 diff --git a/scripts/ci/control/build_blue.sh b/scripts/ci/control/build_blue.sh index 9880e917..c071fec4 100755 --- a/scripts/ci/control/build_blue.sh +++ b/scripts/ci/control/build_blue.sh @@ -266,28 +266,21 @@ run_case top-level-ctest run_ctest "${SOURCE_DIR}/build-ci" run_logged_step "comm-build" bash "${SOURCE_DIR}/tests/comm/build.sh" run_case comm-log "${SOURCE_DIR}/tests/comm/install/bin/test_tilexr_log" run_case comm-spdlog-compile "${SOURCE_DIR}/tests/comm/install/bin/test_tilexr_log_spdlog_compile" -run_case comm-source-guards "${SOURCE_DIR}/tests/comm/install/bin/test_tilexr_source_guards" run_logged_step "udma-build" env BUILD_TILEXR_UDMA_DEMO=OFF \ bash "${SOURCE_DIR}/tests/udma/build.sh" run_case udma-transport-layout "${SOURCE_DIR}/tests/udma/install/bin/test_tilexr_udma_transport_layout" run_case udma-registry "${SOURCE_DIR}/tests/udma/install/bin/test_tilexr_udma_registry" -run_case udma-demo-sources "${SOURCE_DIR}/tests/udma/install/bin/test_tilexr_udma_demo_sources" -run_case udma-source-guard "${SOURCE_DIR}/tests/udma/install/bin/test_tilexr_udma_source_guard" run_logged_step "sdma-build" bash \ "${SOURCE_DIR}/tests/sdma/build.sh" "${ASCEND_HOME_PATH}" run_case sdma-metadata "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_metadata" run_case sdma-api-invalid "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_api_invalid" run_case sdma-transport-disabled "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_transport_disabled" -run_case sdma-comm-wiring "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_comm_wiring" -run_case sdma-source-guard "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_source_guard" run_case sdma-header-compile "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_header_compile" run_logged_step "ep-build" bash "${SOURCE_DIR}/tests/ep/build.sh" full run_case ep-layout "${SOURCE_DIR}/tests/ep/install/bin/test_tilexr_ep_layout" -run_case ep-api-sources "${SOURCE_DIR}/tests/ep/install/bin/test_tilexr_ep_api_sources" -run_case ep-kernel-sources "${SOURCE_DIR}/tests/ep/install/bin/test_tilexr_ep_kernel_sources" run_case ep-host-validation "${SOURCE_DIR}/tests/ep/install/bin/test_tilexr_ep_host_validation" run_logged_step "memory-configure" cmake \ @@ -296,8 +289,8 @@ run_logged_step "memory-configure" cmake \ -DBUILD_TILEXR_MEMORY_DEMO=ON \ -DTILEXR_MEMORY_DEMO_SOC_TYPE=Ascend910B run_logged_step "memory-build" cmake \ - --build "${SOURCE_DIR}/tests/memory/build" --target install -j"${BUILD_JOBS}" -run_case memory-demo-sources "${SOURCE_DIR}/tests/memory/install/bin/test_tilexr_memory_demo_sources" + --build "${SOURCE_DIR}/tests/memory/build" -j"${BUILD_JOBS}" +run_logged_step "memory-install" cmake --install "${SOURCE_DIR}/tests/memory/build" # Dedicated suite builds reuse the root install prefix with optional components # disabled. Restore the complete top-level install before validating outputs. diff --git a/scripts/ci/host_checks.sh b/scripts/ci/host_checks.sh index 5c5816b2..3815cea1 100755 --- a/scripts/ci/host_checks.sh +++ b/scripts/ci/host_checks.sh @@ -159,10 +159,9 @@ run_comm_host_tests() { cmake --build "${BUILD_ROOT}/comm" --target install -j"$(nproc)" "${BUILD_ROOT}/comm-install/bin/test_tilexr_log" "${BUILD_ROOT}/comm-install/bin/test_tilexr_log_spdlog_compile" - "${BUILD_ROOT}/comm-install/bin/test_tilexr_source_guards" } -run_ep_source_tests() { +run_ep_host_tests() { cmake -S "${ROOT_DIR}/tests/ep" -B "${BUILD_ROOT}/ep" \ -DBUILD_TILEXR_EP_DEMO=OFF cmake --build "${BUILD_ROOT}/ep" -j"$(nproc)" @@ -231,12 +230,10 @@ trap finalize_host_checks EXIT run_and_accumulate_case shell-syntax check_tracked_shell_syntax run_and_accumulate_case ci-ctest run_ci_ctest run_and_accumulate_case comm-host run_comm_host_tests -run_and_accumulate_case ep-source-only run_ep_source_tests +run_and_accumulate_case ep-host run_ep_host_tests run_and_accumulate_case data-as-flag run_data_as_flag_tests run_and_accumulate_case collectives-vllm-patch \ python3 -m pytest -q "${ROOT_DIR}/tests/collectives/unit/test_vllm_collectives_patch.py" -run_and_accumulate_case collectives-vllm-integration-sources \ - python3 -m pytest -q "${ROOT_DIR}/tests/collectives/unit/test_vllm_collectives_integration_sources.py" run_and_accumulate_case collectives-profile-report \ python3 "${ROOT_DIR}/tests/collectives/unit/test_collective_profile_report.py" diff --git a/scripts/ci/provision/common.sh b/scripts/ci/provision/common.sh index c53b2466..7ab1990d 100755 --- a/scripts/ci/provision/common.sh +++ b/scripts/ci/provision/common.sh @@ -8,7 +8,7 @@ CI_GROUP=HwHiAiUser CI_HOME=/home/tilexr-ci CANN_HOME=/home/tilexr-ci/toolchains/cann/9.1.0 CANN_OWNER=root -CONTROL_VERSION=v4 +CONTROL_VERSION=v5 CONTROL_HOME="${CI_HOME}/control/${CONTROL_VERSION}" RUNNER_HOME=/home/tilexr-ci/actions-runner GITHUB_PROXY=http://127.0.0.1:3128 diff --git a/tests/checker/CMakeLists.txt b/tests/checker/CMakeLists.txt index e02d34c6..f02547ce 100644 --- a/tests/checker/CMakeLists.txt +++ b/tests/checker/CMakeLists.txt @@ -15,10 +15,6 @@ function(tilexr_checker_shim_test name) ) endfunction() -tilexr_checker_test(test_tilexr_checker_sources - unit/test_checker_sources.cpp -) - tilexr_checker_test(test_tilexr_checker_world unit/test_checker_world.cpp ) @@ -63,10 +59,6 @@ tilexr_checker_test(test_tilexr_checker_trace_schedule unit/test_checker_trace_schedule.cpp ) -tilexr_checker_test(test_tilexr_checker_source_preservation - unit/test_checker_source_preservation.cpp -) - tilexr_checker_test(test_tilexr_checker_ep_dispatch_oracle_probe unit/test_checker_ep_dispatch_oracle_probe.cpp ) @@ -97,7 +89,6 @@ target_compile_features(test_tilexr_checker_allreduce_big_data_header_probe PRIV add_custom_target(tilexr_checker_all DEPENDS tilexr_checker - test_tilexr_checker_sources test_tilexr_checker_world test_tilexr_checker_oracles test_tilexr_checker_shim_events @@ -109,7 +100,6 @@ add_custom_target(tilexr_checker_all test_tilexr_checker_trace_bundle_verifier test_tilexr_checker_trace_materializer test_tilexr_checker_trace_schedule - test_tilexr_checker_source_preservation test_tilexr_checker_ep_dispatch_oracle_probe test_tilexr_checker_header_probe test_tilexr_checker_allgather_hdb_header_probe diff --git a/tests/checker/unit/test_checker_shim_events.cpp b/tests/checker/unit/test_checker_shim_events.cpp index a41161b7..26ba9b1c 100644 --- a/tests/checker/unit/test_checker_shim_events.cpp +++ b/tests/checker/unit/test_checker_shim_events.cpp @@ -1,8 +1,6 @@ #include #include -#include #include -#include #include #include @@ -25,17 +23,6 @@ void UseCheckerShimTypes() { (void)local; } -std::string SourcePath(const std::string &relative_path) { - return std::string(TILEXR_SOURCE_ROOT) + "/" + relative_path; -} - -std::string ReadFile(const std::string &path) { - std::ifstream input(path.c_str()); - std::ostringstream out; - out << input.rdbuf(); - return out.str(); -} - void ExpectTrue(bool condition, const char *message) { if (!condition) { std::cerr << message << "\n"; @@ -88,14 +75,6 @@ void ExpectContains(const std::string &text, const std::string &needle, } } -void ExpectNotContains(const std::string &text, const std::string &needle, - const char *message) { - if (text.find(needle) != std::string::npos) { - std::cerr << message << ": unexpectedly found " << needle << "\n"; - ++g_failures; - } -} - int32_t ReadInt32(const tilexr::checker::ByteBuffer &buffer, size_t index) { int32_t value = 0; tilexr::checker::CheckerStatus status = buffer.ReadInt32(index, &value); @@ -270,44 +249,6 @@ void TestEventsCaptureServerTopology() { ExpectContains(jsonl, "\"peer_server\":1", "events json includes peer server"); } -void TestCheckerLocalShimIncludePathOnly() { - const std::string checker_cmake = ReadFile(SourcePath("tests/checker/CMakeLists.txt")); - const std::string tools_cmake = ReadFile(SourcePath("tools/checker/CMakeLists.txt")); - - ExpectContains(checker_cmake, "test_tilexr_checker_shim_events", - "checker test target wired"); - ExpectContains(checker_cmake, "tools/checker/shim", "checker shim include path wired"); - ExpectNotContains(tools_cmake, "tools/checker/shim", "checker core should not export shim"); -} - -void TestTraceAdapterKeepsAlgorithmShimsThin() { - const std::string common_shim = - ReadFile(SourcePath("tools/checker/shim/tilexr/checker/collective_trace_shim.h")); - const std::string allreduce_adapter = - ReadFile(SourcePath("tools/checker/shim/tilexr/checker/allreduce_big_data_trace_shim.h")); - const std::string hdb_adapter = - ReadFile(SourcePath("tools/checker/shim/tilexr/checker/allgather_hdb_trace_shim.h")); - - ExpectNotContains(common_shim, "#include \"allreduce_big_data.h\"", - "common trace shim should not include a production algorithm"); - ExpectNotContains(common_shim, "#define CpGM2GMPingPong", - "common trace shim should not own production call macros"); - - ExpectContains(allreduce_adapter, "TILEXR_CHECKER_TRACE_TARGET_HEADER", - "allreduce adapter declares target header"); - ExpectContains(allreduce_adapter, "collective_trace_adapter.h", - "allreduce adapter uses generic trace adapter"); - ExpectNotContains(allreduce_adapter, "#define CpGM2GMPingPong", - "allreduce adapter should not duplicate trace macros"); - - ExpectContains(hdb_adapter, "TILEXR_CHECKER_TRACE_TARGET_HEADER", - "hdb adapter declares target header"); - ExpectContains(hdb_adapter, "collective_trace_adapter.h", - "hdb adapter uses generic trace adapter"); - ExpectNotContains(hdb_adapter, "#define CpGM2GMPingPong", - "hdb adapter should not duplicate trace macros"); -} - void TestAscendCPipePrimitivesRecordTraceEvents() { tilexr::checker::CheckerCase test_case; test_case.op = tilexr::checker::CollectiveOp::kAllGather; @@ -677,8 +618,6 @@ int main() { TestNullWorldAndUnsupportedRoleFail(); TestRecordWriteAndBarrierEvents(); TestEventsCaptureServerTopology(); - TestCheckerLocalShimIncludePathOnly(); - TestTraceAdapterKeepsAlgorithmShimsThin(); TestAscendCPipePrimitivesRecordTraceEvents(); TestTraceRuntimeReportsUnresolvedGmCopyAddress(); TestRawAscendCDataCopyRecordsTraceRuntimeCopy(); diff --git a/tests/checker/unit/test_checker_source_preservation.cpp b/tests/checker/unit/test_checker_source_preservation.cpp deleted file mode 100644 index da33501c..00000000 --- a/tests/checker/unit/test_checker_source_preservation.cpp +++ /dev/null @@ -1,117 +0,0 @@ -#include -#include -#include -#include -#include - -#include -#include - -namespace { - -int g_failures = 0; - -std::string SourcePath(const std::string &relative_path) { - return std::string(TILEXR_SOURCE_ROOT) + "/" + relative_path; -} - -std::string ReadFile(const std::string &path) { - std::ifstream input(path.c_str()); - if (!input) { - std::cerr << "failed to open " << path << "\n"; - ++g_failures; - return std::string(); - } - std::ostringstream out; - out << input.rdbuf(); - return out.str(); -} - -void ExpectNotContains(const std::string &text, const std::string &needle, - const std::string &path) { - if (text.find(needle) != std::string::npos) { - std::cerr << path << " unexpectedly contains " << needle << "\n"; - ++g_failures; - } -} - -bool HasAllowedExtension(const std::string &path) { - const std::vector extensions = { - ".h", - ".hpp", - ".cpp", - ".cc", - ".cce", - }; - for (size_t i = 0; i < extensions.size(); ++i) { - const std::string &extension = extensions[i]; - if (path.size() >= extension.size() && - path.compare(path.size() - extension.size(), extension.size(), extension) == 0) { - return true; - } - } - return false; -} - -void AppendProductionSourceFiles(const std::string &directory, std::vector *files) { - DIR *dir = opendir(directory.c_str()); - if (dir == nullptr) { - std::cerr << "failed to open directory " << directory << "\n"; - ++g_failures; - return; - } - - while (true) { - struct dirent *entry = readdir(dir); - if (entry == nullptr) { - break; - } - const std::string name = entry->d_name; - if (name == "." || name == "..") { - continue; - } - const std::string path = directory + "/" + name; - struct stat st; - if (stat(path.c_str(), &st) != 0) { - std::cerr << "failed to stat " << path << "\n"; - ++g_failures; - continue; - } - if (S_ISDIR(st.st_mode)) { - AppendProductionSourceFiles(path, files); - } else if (S_ISREG(st.st_mode) && HasAllowedExtension(path)) { - files->push_back(path); - } - } - closedir(dir); -} - -} // namespace - -int main() { - std::vector production_files; - AppendProductionSourceFiles(SourcePath("src/collectives"), &production_files); - AppendProductionSourceFiles(SourcePath("src/ep"), &production_files); - - if (production_files.empty()) { - std::cerr << "no production source files scanned\n"; - ++g_failures; - } - - const std::vector rejected_tokens = { - "TILEXR_BUILD_CHECKER", - "TILEXR_CHECKER", - "checker::", - "tools/checker", - }; - - for (size_t file_index = 0; file_index < production_files.size(); ++file_index) { - const std::string path = production_files[file_index]; - const std::string text = ReadFile(path); - for (size_t token_index = 0; token_index < rejected_tokens.size(); ++token_index) { - ExpectNotContains(text, rejected_tokens[token_index], path); - } - } - - return g_failures == 0 ? 0 : 1; -} diff --git a/tests/checker/unit/test_checker_sources.cpp b/tests/checker/unit/test_checker_sources.cpp deleted file mode 100644 index e85f80d4..00000000 --- a/tests/checker/unit/test_checker_sources.cpp +++ /dev/null @@ -1,35 +0,0 @@ -#include -#include -#include -#include - -namespace { -int g_failures = 0; - -std::string SourcePath(const std::string &relative_path) { - return std::string(TILEXR_SOURCE_ROOT) + "/" + relative_path; -} - -std::string ReadFile(const std::string &path) { - std::ifstream input(path.c_str()); - std::ostringstream out; - out << input.rdbuf(); - return out.str(); -} - -void ExpectContains(const std::string &text, const std::string &needle, - const std::string &path) { - if (text.find(needle) == std::string::npos) { - std::cerr << path << " missing " << needle << "\n"; - ++g_failures; - } -} -} // namespace - -int main() { - const std::string root_cmake = ReadFile(SourcePath("CMakeLists.txt")); - ExpectContains(root_cmake, "TILEXR_BUILD_CHECKER", "CMakeLists.txt"); - const std::string checker_cmake = ReadFile(SourcePath("tools/checker/CMakeLists.txt")); - ExpectContains(checker_cmake, "tilexr-checker-core", "tools/checker/CMakeLists.txt"); - return g_failures == 0 ? 0 : 1; -} diff --git a/tests/checker/unit/test_checker_trace_adapter_generator.cpp b/tests/checker/unit/test_checker_trace_adapter_generator.cpp index d65588e0..db6ff768 100644 --- a/tests/checker/unit/test_checker_trace_adapter_generator.cpp +++ b/tests/checker/unit/test_checker_trace_adapter_generator.cpp @@ -1054,7 +1054,7 @@ void TestRenderTraceOnboardingPlanDocumentsExecutorWork() { ExpectContains(output, "git diff -- src/collectives src/ep", "onboarding plan source preservation check"); ExpectContains(output, "Do not edit production sources", - "onboarding plan production source guard"); + "onboarding plan production source boundary"); } void TestRenderTraceOnboardingPlanIncludesManualReviewActions() { diff --git a/tests/ci/test_control_contract.py b/tests/ci/test_control_contract.py index adec07f3..0eb89301 100644 --- a/tests/ci/test_control_contract.py +++ b/tests/ci/test_control_contract.py @@ -74,21 +74,13 @@ def test_build_manifest_enables_and_runs_all_non_hardware_coverage(self): '(cd "$1" && ctest --output-on-failure)', "test_tilexr_log", "test_tilexr_log_spdlog_compile", - "test_tilexr_source_guards", "test_tilexr_udma_transport_layout", "test_tilexr_udma_registry", - "test_tilexr_udma_demo_sources", - "test_tilexr_udma_source_guard", "test_tilexr_sdma_metadata", "test_tilexr_sdma_api_invalid", "test_tilexr_sdma_transport_disabled", - "test_tilexr_sdma_comm_wiring", - "test_tilexr_sdma_source_guard", "test_tilexr_sdma_header_compile", - "test_tilexr_memory_demo_sources", "test_tilexr_ep_layout", - "test_tilexr_ep_api_sources", - "test_tilexr_ep_kernel_sources", "test_tilexr_ep_host_validation", "libascend_hal.so", "readelf", @@ -128,14 +120,14 @@ def test_live_runner_requires_python_pidfd_support(self): ]: self.assertIn(token, text) - def test_sealed_controller_v4_is_consistent_across_provisioning(self): - self.assertEqual("v4\n", self.read("scripts/ci/control/VERSION")) + def test_sealed_controller_v5_is_consistent_across_provisioning(self): + self.assertEqual("v5\n", self.read("scripts/ci/control/VERSION")) common = self.read("scripts/ci/provision/common.sh") control = self.read("scripts/ci/provision/control.sh") verify = self.read("scripts/ci/provision/verify.sh") workflow = self.read(".github/workflows/npu-ci.yml") - self.assertIn("CONTROL_VERSION=v4", common) + self.assertIn("CONTROL_VERSION=v5", common) self.assertIn( 'CONTROL_HOME="${CI_HOME}/control/${CONTROL_VERSION}"', common ) @@ -481,7 +473,7 @@ def test_build_manifest_resets_cases_and_records_each_safe_test(self): text = self.read("scripts/ci/control/build_blue.sh") self.assertIn(': > "${CASES_FILE}"', text) self.assertIn("run_case()", text) - self.assertGreaterEqual(text.count("run_case "), 19) + self.assertGreaterEqual(text.count("run_case "), 11) self.assertIn("printf '%s\\t%s\\t%s\\t%s\\n'", text) def test_build_manifest_captures_configuration_and_build_logs(self): @@ -530,12 +522,10 @@ def test_host_checks_run_complete_fast_coverage(self): 'cmake -S "${ROOT_DIR}/tests/comm"', "test_tilexr_log", "test_tilexr_log_spdlog_compile", - "test_tilexr_source_guards", 'cmake -S "${ROOT_DIR}/tests/ep"', "-DBUILD_TILEXR_EP_DEMO=OFF", 'cmake -S "${ROOT_DIR}/tests/data_as_flag"', "test_vllm_collectives_patch.py", - "test_vllm_collectives_integration_sources.py", "test_collective_profile_report.py", "CASE_NAMES=()", "run_case()", @@ -552,10 +542,9 @@ def test_host_checks_run_complete_fast_coverage(self): "shell-syntax", "ci-ctest", "comm-host", - "ep-source-only", + "ep-host", "data-as-flag", "collectives-vllm-patch", - "collectives-vllm-integration-sources", "collectives-profile-report", ]: self.assertIn("run_and_accumulate_case {} ".format(suite), text) @@ -955,8 +944,7 @@ def test_host_checks_runs_all_cases_after_an_intermediate_failure(self): " printf '\\n' > \"$build_dir/Testing/tag/Test.xml\"\n" "fi\n" "mkdir -p \"${TILEXR_TEST_COMM_BIN:?}\"\n" - "for name in test_tilexr_log test_tilexr_log_spdlog_compile " - "test_tilexr_source_guards; do\n" + "for name in test_tilexr_log test_tilexr_log_spdlog_compile; do\n" " printf '#!/bin/bash\\nexit 0\\n' > " "\"${TILEXR_TEST_COMM_BIN}/$name\"\n" " chmod 755 \"${TILEXR_TEST_COMM_BIN}/$name\"\n" @@ -1017,29 +1005,28 @@ def test_host_checks_runs_all_cases_after_an_intermediate_failure(self): line for line in result.stdout.splitlines() if line.startswith("[host-check] ") ] - self.assertEqual(8, len(case_lines)) + self.assertEqual(7, len(case_lines)) expected = [ "shell-syntax", "ci-ctest", "comm-host", - "ep-source-only", + "ep-host", "data-as-flag", "collectives-vllm-patch", - "collectives-vllm-integration-sources", "collectives-profile-report", ] self.assertEqual( expected, [line.split(":", 1)[0][len("[host-check] "):] for line in case_lines], ) - self.assertIn("[host-check] ep-source-only: FAIL (exit=23", result.stdout) + self.assertIn("[host-check] ep-host: FAIL (exit=23", result.stdout) self.assertIn( "[host-check] collectives-profile-report: FAIL (exit=31", result.stdout, ) - self.assertIn("Host cases: 8 total, 6 passed, 2 failed", result.stdout) + self.assertIn("Host cases: 7 total, 5 passed, 2 failed", result.stdout) self.assertIn( - "Failed cases: ep-source-only, collectives-profile-report", + "Failed cases: ep-host, collectives-profile-report", result.stdout, ) diff --git a/tests/collectives/CMakeLists.txt b/tests/collectives/CMakeLists.txt index d4feeb0e..d9061910 100644 --- a/tests/collectives/CMakeLists.txt +++ b/tests/collectives/CMakeLists.txt @@ -62,30 +62,6 @@ if(TARGET tilexr-collectives) set(TILEXR_COLLECTIVES_IN_TREE ON) endif() -add_executable(test_tilexr_collectives_api - unit/test_tilexr_collectives_api.cpp -) - -target_compile_definitions(test_tilexr_collectives_api PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - -add_executable(test_tilexr_collectives_kernel_ownership - unit/test_tilexr_collectives_kernel_ownership.cpp -) - -target_compile_definitions(test_tilexr_collectives_kernel_ownership PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - -add_executable(test_tilexr_collectives_tools_sources - unit/test_tilexr_collectives_tools_sources.cpp -) - -target_compile_definitions(test_tilexr_collectives_tools_sources PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - add_executable(test_tilexr_collectives_header_compile unit/test_tilexr_collectives_header_compile.cpp ) @@ -308,9 +284,6 @@ foreach(_tilexr_collectives_link_target TileXRAddAscendRpathLink(${_tilexr_collectives_link_target}) endforeach() -add_test(NAME test_tilexr_collectives_api COMMAND test_tilexr_collectives_api) -add_test(NAME test_tilexr_collectives_kernel_ownership COMMAND test_tilexr_collectives_kernel_ownership) -add_test(NAME test_tilexr_collectives_tools_sources COMMAND test_tilexr_collectives_tools_sources) add_test(NAME test_tilexr_collectives_header_compile COMMAND test_tilexr_collectives_header_compile) add_test(NAME test_tilexr_perf_trace_layout COMMAND test_tilexr_perf_trace_layout) add_test(NAME test_collective_perf_report COMMAND test_collective_perf_report) @@ -326,17 +299,9 @@ if(Python3_Interpreter_FOUND) COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/unit/test_collective_profile_report.py ) - add_test( - NAME test_vllm_collectives_integration_sources - COMMAND ${Python3_EXECUTABLE} - ${CMAKE_CURRENT_SOURCE_DIR}/unit/test_vllm_collectives_integration_sources.py - ) endif() install(TARGETS - test_tilexr_collectives_api - test_tilexr_collectives_kernel_ownership - test_tilexr_collectives_tools_sources test_tilexr_collectives_header_compile test_tilexr_perf_trace_layout test_collective_perf_report diff --git a/tests/collectives/unit/test_tilexr_collectives_api.cpp b/tests/collectives/unit/test_tilexr_collectives_api.cpp deleted file mode 100644 index 877e6dcf..00000000 --- a/tests/collectives/unit/test_tilexr_collectives_api.cpp +++ /dev/null @@ -1,434 +0,0 @@ -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -#define CHECK_TRUE(expr) \ - do { \ - if (!(expr)) { \ - std::cerr << "CHECK_TRUE failed at line " << __LINE__ << ": " #expr << std::endl; \ - ++g_failures; \ - } \ - } while (0) - -std::string RepoPath(const std::string& path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -std::string ReadFile(const std::string& path) -{ - const std::string fullPath = RepoPath(path); - std::ifstream input(fullPath.c_str()); - if (!input.is_open()) { - std::cerr << "failed to open " << fullPath << std::endl; - ++g_failures; - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -void CheckFileDoesNotExist(const std::string& path) -{ - const std::string fullPath = RepoPath(path); - std::ifstream input(fullPath.c_str()); - if (input.is_open()) { - std::cerr << "expected " << fullPath << " not to exist" << std::endl; - ++g_failures; - } -} - -void CheckContains(const std::string& path, const std::string& text, const std::string& needle) -{ - const auto pos = text.find(needle); - if (pos == std::string::npos) { - std::cerr << "expected " << path << " to contain: " << needle << std::endl; - ++g_failures; - } -} - -void CheckDoesNotContain(const std::string& path, const std::string& text, const std::string& needle) -{ - const auto pos = text.find(needle); - if (pos != std::string::npos) { - std::cerr << "expected " << path << " not to contain: " << needle - << " at byte " << pos << std::endl; - ++g_failures; - } -} - -void TestCollectivesHeaderDeclaresPublicApis() -{ - const std::string path = "src/include/tilexr_collectives.h"; - const auto text = ReadFile(path); - CheckContains(path, text, "#ifdef __cplusplus"); - CheckContains(path, text, "extern \"C\""); - CheckContains(path, text, "int TileXRAllGather(void *sendBuf, void *recvBuf, int64_t sendCount,"); - CheckContains(path, text, "int TileXRAllToAll(void *sendBuf, void *recvBuf, int64_t sendCount,"); - CheckContains(path, text, "int TileXRAllReduce(void *sendBuf, void *recvBuf, int64_t count,"); - CheckContains(path, text, "TileXR::TileXRReduceOp op,"); - CheckContains(path, text, "int TileXRReduceScatter(void *sendBuf, void *recvBuf, int64_t recvCount,"); - CheckContains(path, text, "int TileXRBroadcast(void *buf, int64_t count,"); -} - -void TestCoreApiHeaderDoesNotDeclareCollectives() -{ - const std::string path = "src/include/tilexr_api.h"; - const auto text = ReadFile(path); - CheckContains(path, text, "int TileXRCommNextMagic(TileXRCommPtr comm, int64_t *magic);"); - CheckDoesNotContain(path, text, "TileXRAllGather"); - CheckDoesNotContain(path, text, "TileXRAllToAll"); - CheckDoesNotContain(path, text, "TileXRAllReduce"); - CheckDoesNotContain(path, text, "TileXRReduceScatter"); - CheckDoesNotContain(path, text, "TileXRBroadcast"); -} - -void TestCollectivesHostUsesOnlyPublicCommExtensionApi() -{ - const std::string paths[] = { - "src/collectives/host/tilexr_collectives.cpp", - "src/collectives/host/collective_launcher.h", - "src/collectives/host/collective_launcher.cpp", - "src/collectives/host/collective_utils.h", - "src/collectives/host/collective_utils.cpp", - "src/collectives/host/collective_kernel.h", - "src/collectives/host/collective_kernel.cpp", - }; - - bool sawCommArgsHost = false; - bool sawCommArgsDev = false; - bool sawNextMagic = false; - for (const auto& path : paths) { - const auto text = ReadFile(path); - CheckDoesNotContain(path, text, "tilexr_comm.h"); - sawCommArgsHost = sawCommArgsHost || text.find("TileXRGetCommArgsHost") != std::string::npos; - sawCommArgsDev = sawCommArgsDev || text.find("TileXRGetCommArgsDev") != std::string::npos; - sawNextMagic = sawNextMagic || text.find("TileXRCommNextMagic") != std::string::npos; - } - - CHECK_TRUE(sawCommArgsHost); - CHECK_TRUE(sawCommArgsDev); - CHECK_TRUE(sawNextMagic); - - const std::string launcherPath = "src/collectives/host/collective_launcher.cpp"; - CheckDoesNotContain(launcherPath, ReadFile(launcherPath), "TileXRCommNextMagic"); - - const std::string kernelPath = "src/collectives/host/collective_kernel.cpp"; - CheckContains(kernelPath, ReadFile(kernelPath), "TileXRCommNextMagic"); -} - -void TestCollectivesHostOwnsCollectiveLaunchHelpers() -{ - const std::string utilsHeaderPath = "src/collectives/host/collective_utils.h"; - const auto utilsHeader = ReadFile(utilsHeaderPath); - CheckContains(utilsHeaderPath, utilsHeader, "bool IsSupportedDataType(TileXR::TileXRDataType dataType);"); - CheckContains(utilsHeaderPath, utilsHeader, "bool IsSupportedReductionDataType(TileXR::TileXRDataType dataType);"); - CheckContains(utilsHeaderPath, utilsHeader, "int64_t CountToBytes(int64_t count, TileXR::TileXRDataType dataType);"); - CheckContains(utilsHeaderPath, utilsHeader, "uint32_t GetAllGatherBlockNum(const TileXR::CommArgs &commArgs, int64_t dataSize);"); - CheckContains(utilsHeaderPath, utilsHeader, "uint32_t GetAllToAllBlockNum(const TileXR::CommArgs &commArgs, int64_t dataSize);"); - - const std::string kernelHeaderPath = "src/collectives/host/collective_kernel.h"; - const auto kernelHeader = ReadFile(kernelHeaderPath); - CheckContains(kernelHeaderPath, kernelHeader, "namespace TileXRCollectives"); - CheckContains(kernelHeaderPath, kernelHeader, "namespace Host"); - CheckContains(kernelHeaderPath, kernelHeader, "struct AscendCCLKernelArgs"); - CheckContains(kernelHeaderPath, kernelHeader, "int LaunchCollectiveKernel(TileXRCommPtr comm, TileXR::TileXRType type,"); -} - -void TestBroadcastLaunchUsesByteCount() -{ - const std::string path = "src/collectives/host/tilexr_collectives.cpp"; - const auto text = ReadFile(path); - CheckContains(path, text, "const int64_t bytes = TileXRCollectives::Host::CountToBytes(count, dataType);"); - CheckContains(path, text, - "buf, buf, bytes, dataType, blockDim, stream,\n" - " TileXRCollectives::Host::CollectiveLaunchAttrs { 0, root }"); -} - -void TestCommBuildDoesNotReferenceCollectives() -{ - const std::string path = "src/comm/CMakeLists.txt"; - const auto text = ReadFile(path); - CheckDoesNotContain(path, text, "src/collectives"); - CheckDoesNotContain(path, text, "tilexr_collectives"); -} - -void TestCommInternalDoesNotContainCollectiveRegistration() -{ - const std::string paths[] = { - "src/comm/tilexr_internal.h", - "src/comm/tilexr_internal.cpp", - "src/comm/tilexr_comm.cpp", - }; - const std::string forbidden[] = { - "TILEXR_CCE_BIN_STR", - "RegistCCL", - "RegistCoCKernel", - "RegistKernel", - "LoadMTE", - "LaunchCollective", - "AscendCCLKernelArgs", - "rtKernelLaunch", - }; - - for (const auto& path : paths) { - const auto text = ReadFile(path); - for (const auto& needle : forbidden) { - CheckDoesNotContain(path, text, needle); - } - } - - CheckFileDoesNotExist("src/comm/ccl_kernel_args.h"); -} - -void TestCommBuildInstallsPublicHeadersAndKeepsLinksPrivate() -{ - const std::string path = "src/comm/CMakeLists.txt"; - const auto text = ReadFile(path); - CheckContains(path, text, "include(GNUInstallDirs)"); - CheckContains(path, text, "target_link_directories(tile-comm\n PRIVATE"); - CheckContains(path, text, "target_link_libraries(tile-comm\n PRIVATE"); - CheckDoesNotContain(path, text, "target_link_libraries(tile-comm ascendcl"); - CheckContains(path, text, "LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR}"); - CheckContains(path, text, "DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}"); - CheckContains(path, text, "${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_api.h"); - CheckContains(path, text, "${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_types.h"); - CheckContains(path, text, "${CMAKE_CURRENT_SOURCE_DIR}/../include/comm_args.h"); - CheckContains(path, text, "${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_sync.h"); - CheckContains(path, text, "${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_udma.h"); - CheckContains(path, text, "${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_udma_reg.h"); - CheckContains(path, text, "${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_udma_types.h"); - CheckDoesNotContain(path, text, "${CMAKE_SOURCE_DIR}/src/include/"); - CheckDoesNotContain(path, text, "${CMAKE_SOURCE_DIR}/src/include/tilexr_collectives.h"); - CheckDoesNotContain(path, text, "FILES_MATCHING PATTERN \"*.h\""); -} - -void TestCollectivesBuildDefinesSeparateSharedLibrary() -{ - const std::string path = "src/collectives/CMakeLists.txt"; - const auto text = ReadFile(path); - CheckContains(path, text, "include(GNUInstallDirs)"); - CheckContains(path, text, "host/collective_launcher.cpp"); - CheckContains(path, text, "host/collective_utils.cpp"); - CheckContains(path, text, "host/collective_kernel.cpp"); - CheckContains(path, text, "add_subdirectory(kernels)"); - CheckContains(path, text, "tilexr_collectives_kernel_embed.cpp"); - CheckContains(path, text, "tilexr_collectives_op"); - CheckContains(path, text, "add_library(tilexr-collectives SHARED"); - CheckContains(path, text, "${ASCEND_DRIVER_PATH}/kernel/inc\n PRIVATE"); - CheckContains(path, text, "target_link_libraries(tilexr-collectives\n PUBLIC\n tile-comm\n PRIVATE"); - CheckContains(path, text, "target_link_directories(tilexr-collectives\n PRIVATE"); - CheckDoesNotContain(path, text, "target_link_libraries(tilexr-collectives tile-comm"); - CheckContains(path, text, "LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR}"); - CheckContains(path, text, "${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_collectives.h"); - CheckContains(path, text, "DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}"); - CheckDoesNotContain(path, text, "${CMAKE_SOURCE_DIR}/src/include/"); - CheckDoesNotContain(path, text, "tilexr_api.h"); - CheckDoesNotContain(path, text, "comm_args.h"); - CheckDoesNotContain(path, text, "${CMAKE_INSTALL_PREFIX}/lib"); - CheckDoesNotContain(path, text, "${CMAKE_INSTALL_PREFIX}/include"); -} - -void TestCollectivesKernelBuildSupportsAscend950Arch() -{ - const std::string path = "src/collectives/kernels/CMakeLists.txt"; - const auto text = ReadFile(path); - - CheckContains(path, text, "TILEXR_COLLECTIVES_SOC_TYPE"); - CheckContains(path, text, "TILEXR_SOC_NAME"); - CheckContains(path, text, "Ascend950"); - CheckContains(path, text, "MATCHES \"^Ascend950\""); - CheckContains(path, text, "dav-c310-vec"); - CheckContains(path, text, "dav-c220-vec"); - CheckContains(path, text, "--cce-aicore-arch=${TILEXR_COLLECTIVES_AICORE_ARCH}"); - - const std::string collectivesPath = "src/collectives/CMakeLists.txt"; - const auto collectivesText = ReadFile(collectivesPath); - CheckContains(collectivesPath, collectivesText, "TILEXR_COLLECTIVES_C310_ATOMIC_LIMITS"); - - const std::string utilsPath = "src/collectives/host/collective_utils.cpp"; - const auto utilsText = ReadFile(utilsPath); - CheckContains(utilsPath, utilsText, "TILEXR_COLLECTIVES_C310_ATOMIC_LIMITS"); - CheckContains(utilsPath, utilsText, "TILEXR_DATA_TYPE_INT64"); - - const std::string apiPath = "src/collectives/host/tilexr_collectives.cpp"; - const auto apiText = ReadFile(apiPath); - CheckContains(apiPath, apiText, "IsSupportedReductionDataType(dataType)"); -} - -void TestCollectivesKernelSourcesAllowAscend950Macros() -{ - const std::string kernelTuPath = "src/collectives/kernels/tilexr_lccl_op.cpp"; - const auto kernelTu = ReadFile(kernelTuPath); - CheckContains(kernelTuPath, kernelTu, "__DAV_C220_VEC__"); - CheckContains(kernelTuPath, kernelTu, "__DAV_C310_VEC__"); - CheckContains(kernelTuPath, kernelTu, "LCCL_TYPE_AIV_FUNC(LCCL_ALLGATHER_FUNC_AUTO_DEF)"); - - const std::string lcclOpPath = "src/collectives/kernels/lccl_op.h"; - const auto lcclOp = ReadFile(lcclOpPath); - CheckContains(lcclOpPath, lcclOp, "__DAV_C220_VEC__"); - CheckContains(lcclOpPath, lcclOp, "__DAV_C310_VEC__"); - CheckContains(lcclOpPath, lcclOp, "TileXRAllGather_##type##suffix"); - - const std::string collectivesCcePath = "src/collectives/kernels/kernels/collectives.cce"; - const auto collectivesCce = ReadFile(collectivesCcePath); - CheckContains(collectivesCcePath, collectivesCce, "__DAV_C220_VEC__"); - CheckContains(collectivesCcePath, collectivesCce, "__DAV_C310_VEC__"); - CheckContains(collectivesCcePath, collectivesCce, "#define __aicore__"); - CheckContains(collectivesCcePath, collectivesCce, "#if defined(__DAV_C220_VEC__) || defined(__DAV_C310_VEC__)"); - - const std::string collectivesHeaderPath = "src/collectives/kernels/collectives.h"; - const auto collectivesHeader = ReadFile(collectivesHeaderPath); - CheckContains(collectivesHeaderPath, collectivesHeader, "#if defined(__DAV_C220_VEC__) || defined(__DAV_C310_VEC__)"); - - const std::string dataCopyPath = "src/collectives/kernels/datacopy_gm2gm.h"; - const auto dataCopy = ReadFile(dataCopyPath); - CheckContains(dataCopyPath, dataCopy, "#if defined(__DAV_C220_VEC__) || defined(__DAV_C310_VEC__)"); - CheckContains(dataCopyPath, dataCopy, "TileXRAtomicTypeSupported"); - CheckContains(dataCopyPath, dataCopy, "std::is_same_v"); - CheckContains(dataCopyPath, dataCopy, "if constexpr (TileXRAtomicTypeSupported::value)"); -} - -void TestRootBuildRegistersCollectivesTests() -{ - const std::string path = "CMakeLists.txt"; - const auto text = ReadFile(path); - CheckContains(path, text, "option(TILEXR_BUILD_TESTS \"Build TileXR tests\" OFF)"); - CheckContains(path, text, "include(CTest)"); - CheckContains(path, text, "include(GNUInstallDirs)"); - CheckContains(path, text, "if(TILEXR_BUILD_COLLECTIVES)"); - CheckContains(path, text, "add_subdirectory(src/collectives)"); - CheckContains(path, text, "if(BUILD_TESTING OR TILEXR_BUILD_TESTS)"); - CheckContains(path, text, "add_subdirectory(tests/collectives)"); - CheckContains(path, text, "tilexr_collectives_install_smoke.cmake"); - CheckContains(path, text, "tests/collectives/cmake/install_prefix_smoke.cmake.in"); - CheckContains(path, text, "add_test(NAME tilexr_collectives_install_smoke"); - CheckContains(path, text, "${CMAKE_COMMAND} -P"); - CheckContains(path, text, "tilexr_collectives_install_smoke"); -} - -void TestCollectivesTestBuildSupportsInTreeAndInstallPrefixModes() -{ - const std::string path = "tests/collectives/CMakeLists.txt"; - const auto text = ReadFile(path); - CheckContains(path, text, "include(GNUInstallDirs)"); - CheckContains(path, text, "enable_testing()"); - CheckContains(path, text, "set(TILEXR_COLLECTIVES_IN_TREE OFF)"); - CheckContains(path, text, "if(TARGET tilexr-collectives)"); - CheckContains(path, text, "set(TILEXR_COLLECTIVES_IN_TREE ON)"); - CheckContains(path, text, "if(TILEXR_COLLECTIVES_IN_TREE)"); - CheckContains(path, text, "set(TILEXR_COLLECTIVES_TEST_TARGET tilexr-collectives)"); - CheckContains(path, text, "else()"); - CheckContains(path, text, "set(TILEXR_INSTALL_PREFIX \"${TILEXR_ROOT}/install\" CACHE PATH"); - CheckContains(path, text, "set(TILEXR_INSTALL_LIBDIR \"${CMAKE_INSTALL_LIBDIR}\" CACHE STRING"); - CheckContains(path, text, "set(TILEXR_INSTALL_INCLUDEDIR \"${CMAKE_INSTALL_INCLUDEDIR}\" CACHE STRING"); - CheckContains(path, text, "if(IS_ABSOLUTE \"${TILEXR_INSTALL_LIBDIR}\")"); - CheckContains(path, text, "set(TILEXR_INSTALL_LIB_SEARCH_DIR \"${TILEXR_INSTALL_LIBDIR}\")"); - CheckContains(path, text, "set(TILEXR_INSTALL_LIB_SEARCH_DIR \"${TILEXR_INSTALL_PREFIX}/${TILEXR_INSTALL_LIBDIR}\")"); - CheckContains(path, text, "if(IS_ABSOLUTE \"${TILEXR_INSTALL_INCLUDEDIR}\")"); - CheckContains(path, text, "set(TILEXR_INSTALL_INCLUDE_SEARCH_DIR \"${TILEXR_INSTALL_INCLUDEDIR}\")"); - CheckContains(path, text, "set(TILEXR_INSTALL_INCLUDE_SEARCH_DIR \"${TILEXR_INSTALL_PREFIX}/${TILEXR_INSTALL_INCLUDEDIR}\")"); - CheckContains(path, text, "${TILEXR_INSTALL_INCLUDE_SEARCH_DIR}"); - CheckContains(path, text, "${TILEXR_INSTALL_INCLUDE_SEARCH_DIR}/tilexr_collectives.h"); - CheckContains(path, text, "set(ASCEND_HOME_PATH \"${_tilexr_default_ascend_home_path}\" CACHE PATH"); - CheckContains(path, text, "set(ARCH \"${_tilexr_default_arch}\" CACHE STRING"); - CheckContains(path, text, "set(ASCEND_DRIVER_PATH \"${_tilexr_default_ascend_driver_path}\" CACHE PATH"); - CheckContains(path, text, "set(TILEXR_COLLECTIVES_LIB \"\" CACHE FILEPATH"); - CheckContains(path, text, "set(TILEXR_LIB \"\" CACHE FILEPATH"); - CheckContains(path, text, "${TILEXR_INSTALL_LIB_SEARCH_DIR}"); - CheckContains(path, text, "foreach(_tilexr_installed_include_dir"); - CheckContains(path, text, "if(EXISTS \"${_tilexr_installed_include_dir}\")"); - CheckContains(path, text, "list(APPEND TILEXR_INSTALLED_INCLUDE_DIRS \"${_tilexr_installed_include_dir}\")"); - CheckContains(path, text, "add_library(tilexr-comm-installed SHARED IMPORTED)"); - CheckContains(path, text, "add_library(tilexr-collectives-installed SHARED IMPORTED)"); - CheckContains(path, text, "IMPORTED_LOCATION \"${TILEXR_LIB}\""); - CheckContains(path, text, "IMPORTED_LOCATION \"${TILEXR_COLLECTIVES_LIB}\""); - CheckContains(path, text, "INTERFACE_LINK_LIBRARIES tilexr-comm-installed"); - CheckContains(path, text, "set(TILEXR_COLLECTIVES_TEST_TARGET tilexr-collectives-installed)"); - CheckContains(path, text, "endif()"); - CheckContains(path, text, "target_link_libraries(test_tilexr_collectives_header_compile\n ${TILEXR_COLLECTIVES_TEST_TARGET}\n)"); - CheckContains(path, text, "target_link_libraries(test_tilexr_collectives_stub_behavior\n ${TILEXR_COLLECTIVES_TEST_TARGET}\n)"); - CheckContains(path, text, "RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}"); - CheckContains(path, text, "add_executable(test_tilexr_collectives_stub_behavior"); - CheckContains(path, text, "unit/test_tilexr_collectives_stub_behavior.cpp"); - CheckContains(path, text, "add_test(NAME test_tilexr_collectives_api COMMAND test_tilexr_collectives_api)"); - CheckContains(path, text, "add_test(NAME test_tilexr_collectives_header_compile COMMAND test_tilexr_collectives_header_compile)"); - CheckContains(path, text, "add_test(NAME test_tilexr_collectives_stub_behavior COMMAND test_tilexr_collectives_stub_behavior)"); - CheckContains(path, text, "message(FATAL_ERROR"); - CheckDoesNotContain(path, text, "${TILEXR_ROOT}/src/include"); - CheckDoesNotContain(path, text, "${TILEXR_ROOT}/3rdparty"); - CheckDoesNotContain(path, text, "${TILEXR_ROOT}/build"); - CheckDoesNotContain(path, text, "${CMAKE_SOURCE_DIR}/src/include"); - CheckDoesNotContain(path, text, "${CMAKE_INSTALL_PREFIX}/bin"); - CheckDoesNotContain(path, text, "${TILEXR_INSTALL_PREFIX}/lib"); - CheckDoesNotContain(path, text, "target_include_directories(test_tilexr_collectives_header_compile"); - CheckDoesNotContain(path, text, "target_include_directories(test_tilexr_collectives_stub_behavior"); - CheckDoesNotContain(path, text, "${TILEXR_COLLECTIVES_LIB}\n ${TILEXR_LIB}"); - CheckDoesNotContain(path, text, "CACHE PATH \"TileXR install library directory"); - CheckDoesNotContain(path, text, "${TILEXR_INSTALL_PREFIX}/include"); - CheckDoesNotContain(path, text, "/tmp/tilexr-install-split-collectives"); - CheckDoesNotContain(path, text, "/tmp/tilexr-build-split-collectives"); -} - -void TestCollectivesInstallSmokeUsesStandaloneInstallPrefixMode() -{ - const std::string path = "tests/collectives/cmake/install_prefix_smoke.cmake.in"; - const auto text = ReadFile(path); - CheckContains(path, text, "tilexr_collectives_ctest_install"); - CheckContains(path, text, "collectives_install_smoke_build"); - CheckContains(path, text, "cmake --install"); - CheckContains(path, text, "COMMAND \"${CMAKE_COMMAND}\" --install \"${TILEXR_ROOT_BINARY_DIR}\" --prefix \"${TILEXR_SMOKE_INSTALL_PREFIX}\""); - CheckContains(path, text, "COMMAND \"${CMAKE_COMMAND}\" -S \"${TILEXR_COLLECTIVES_TEST_SOURCE_DIR}\" -B \"${TILEXR_SMOKE_BUILD_DIR}\""); - CheckContains(path, text, "-DTILEXR_INSTALL_PREFIX=${TILEXR_SMOKE_INSTALL_PREFIX}"); - CheckContains(path, text, "-DTILEXR_INSTALL_INCLUDEDIR=${TILEXR_SMOKE_INSTALL_INCLUDEDIR}"); - CheckContains(path, text, "COMMAND \"${CMAKE_COMMAND}\" --build \"${TILEXR_SMOKE_BUILD_DIR}\""); - CheckContains(path, text, "COMMAND \"${CMAKE_COMMAND}\" -E env"); - CheckContains(path, text, "\"${CMAKE_CTEST_COMMAND}\" --test-dir \"${TILEXR_SMOKE_BUILD_DIR}\" --output-on-failure"); - CheckContains(path, text, "LD_LIBRARY_PATH=${TILEXR_SMOKE_LD_LIBRARY_PATH}"); - CheckContains(path, text, "find_program(TILEXR_READELF readelf)"); - CheckContains(path, text, "readelf -d"); - CheckContains(path, text, "libtile-comm.so"); - CheckContains(path, text, "TILEXR_SMOKE_COLLECTIVES_LIB"); - CheckContains(path, text, "-DASCEND_HOME_PATH=${TILEXR_SMOKE_ASCEND_HOME_PATH}"); - CheckContains(path, text, "-DARCH=${TILEXR_SMOKE_ARCH}"); - CheckContains(path, text, "-DASCEND_DRIVER_PATH=${TILEXR_SMOKE_ASCEND_DRIVER_PATH}"); - CheckDoesNotContain(path, text, "tilexr-collectives-installed"); - CheckDoesNotContain(path, text, "${TILEXR_ROOT_SOURCE_DIR}/src/collectives"); - CheckDoesNotContain(path, text, "${TILEXR_ROOT_BINARY_DIR}/src/collectives"); -} - -} // namespace - -int main() -{ - TestCollectivesHeaderDeclaresPublicApis(); - TestCoreApiHeaderDoesNotDeclareCollectives(); - TestCollectivesHostUsesOnlyPublicCommExtensionApi(); - TestCollectivesHostOwnsCollectiveLaunchHelpers(); - TestBroadcastLaunchUsesByteCount(); - TestCommBuildDoesNotReferenceCollectives(); - TestCommInternalDoesNotContainCollectiveRegistration(); - TestCommBuildInstallsPublicHeadersAndKeepsLinksPrivate(); - TestCollectivesBuildDefinesSeparateSharedLibrary(); - TestCollectivesKernelBuildSupportsAscend950Arch(); - TestCollectivesKernelSourcesAllowAscend950Macros(); - TestRootBuildRegistersCollectivesTests(); - TestCollectivesTestBuildSupportsInTreeAndInstallPrefixModes(); - TestCollectivesInstallSmokeUsesStandaloneInstallPrefixMode(); - if (g_failures != 0) { - std::cerr << g_failures << " collectives API split checks failed" << std::endl; - return 1; - } - std::cout << "TileXR collectives API split checks passed" << std::endl; - return 0; -} diff --git a/tests/collectives/unit/test_tilexr_collectives_kernel_ownership.cpp b/tests/collectives/unit/test_tilexr_collectives_kernel_ownership.cpp deleted file mode 100644 index 55fd34c9..00000000 --- a/tests/collectives/unit/test_tilexr_collectives_kernel_ownership.cpp +++ /dev/null @@ -1,439 +0,0 @@ -#include - -#include -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -std::string RepoPath(const std::string &path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -std::string ReadFile(const std::string &path) -{ - const std::string fullPath = RepoPath(path); - std::ifstream input(fullPath.c_str()); - if (!input.is_open()) { - std::cerr << "failed to open " << fullPath << std::endl; - ++g_failures; - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -std::string ExtractInitializer(const std::string &path, const std::string &text, const std::string &symbol) -{ - const auto symbolPos = text.find(symbol); - if (symbolPos == std::string::npos) { - std::cerr << "expected " << path << " to contain initializer symbol: " << symbol << std::endl; - ++g_failures; - return {}; - } - - const auto openBrace = text.find('{', symbolPos); - if (openBrace == std::string::npos) { - std::cerr << "expected " << path << " to contain initializer body for: " << symbol << std::endl; - ++g_failures; - return {}; - } - - int depth = 0; - for (std::size_t pos = openBrace; pos < text.size(); ++pos) { - if (text[pos] == '{') { - ++depth; - continue; - } - if (text[pos] == '}') { - --depth; - if (depth == 0) { - return text.substr(openBrace + 1, pos - openBrace - 1); - } - } - } - - std::cerr << "expected " << path << " to contain matching '}' for initializer: " << symbol << std::endl; - ++g_failures; - return {}; -} - -bool DirectoryExists(const std::string &path) -{ - DIR *dir = opendir(RepoPath(path).c_str()); - if (dir == nullptr) { - return false; - } - closedir(dir); - return true; -} - -std::vector CollectFiles(const std::string &path) -{ - std::vector files; - DIR *dir = opendir(RepoPath(path).c_str()); - if (dir == nullptr) { - std::cerr << "failed to open directory " << RepoPath(path) << std::endl; - ++g_failures; - return files; - } - - while (dirent *entry = readdir(dir)) { - const std::string name = entry->d_name; - if (name == "." || name == "..") { - continue; - } - const std::string child = path + "/" + name; - if (entry->d_type == DT_DIR) { - const auto childFiles = CollectFiles(child); - files.insert(files.end(), childFiles.begin(), childFiles.end()); - } else if (entry->d_type == DT_REG) { - files.push_back(child); - } - } - closedir(dir); - return files; -} - -void CheckTrue(bool condition, const std::string &message) -{ - if (!condition) { - std::cerr << message << std::endl; - ++g_failures; - } -} - -void CheckContains(const std::string &path, const std::string &text, const std::string &needle) -{ - if (text.find(needle) == std::string::npos) { - std::cerr << "expected " << path << " to contain: " << needle << std::endl; - ++g_failures; - } -} - -void CheckDoesNotContain(const std::string &path, const std::string &text, const std::string &needle) -{ - const auto pos = text.find(needle); - if (pos != std::string::npos) { - std::cerr << "expected " << path << " not to contain: " << needle - << " at byte " << pos << std::endl; - ++g_failures; - } -} - -void TestCollectivesOwnsCceBuild() -{ - CheckTrue(DirectoryExists("src/collectives/kernels"), - "expected src/collectives/kernels to exist"); - - const std::string collectivesCmakePath = "src/collectives/CMakeLists.txt"; - const auto collectivesCmake = ReadFile(collectivesCmakePath); - CheckContains(collectivesCmakePath, collectivesCmake, "add_subdirectory(kernels)"); - CheckContains(collectivesCmakePath, collectivesCmake, "tilexr_collectives_kernel_embed"); - CheckContains(collectivesCmakePath, collectivesCmake, "tilexr_collectives_op"); - CheckContains(collectivesCmakePath, collectivesCmake, "OBJECT_DEPENDS"); - CheckContains(collectivesCmakePath, collectivesCmake, "${TILEXR_COLLECTIVES_OP}"); - - const std::string kernelsCmakePath = "src/collectives/kernels/CMakeLists.txt"; - const auto kernelsCmake = ReadFile(kernelsCmakePath); - CheckContains(kernelsCmakePath, kernelsCmake, "enable_language(CCE)"); - CheckContains(kernelsCmakePath, kernelsCmake, "tilexr_lccl_op.cpp"); - CheckContains(kernelsCmakePath, kernelsCmake, "TILEXR_COLLECTIVES_KERNEL_DEPENDS"); - CheckContains(kernelsCmakePath, kernelsCmake, "CONFIGURE_DEPENDS"); - CheckContains(kernelsCmakePath, kernelsCmake, "OBJECT_DEPENDS"); - CheckContains(kernelsCmakePath, kernelsCmake, "tilexr_collectives_op.o"); - CheckContains(kernelsCmakePath, kernelsCmake, "tilexr_collectives_op"); - CheckContains(kernelsCmakePath, kernelsCmake, "TILEXR_COLLECTIVES_ENABLE_PROFILING"); - CheckContains(kernelsCmakePath, kernelsCmake, "TILEXR_COLLECTIVES_1OP_BIN_SIZE 10485760"); - CheckDoesNotContain(kernelsCmakePath, kernelsCmake, "src/comm"); -} - -void TestCollectivesKernelSourcesAreScoped() -{ - const std::string kernelTuPath = "src/collectives/kernels/tilexr_lccl_op.cpp"; - const auto kernelTu = ReadFile(kernelTuPath); - CheckContains(kernelTuPath, kernelTu, "LCCL_TYPE_AIV_FUNC(LCCL_ALLGATHER_FUNC_AUTO_DEF)"); - CheckContains(kernelTuPath, kernelTu, "LCCL_TYPE_AIV_FUNC(LCCL_ALL2ALL_FUNC_AUTO_DEF)"); - CheckContains(kernelTuPath, kernelTu, "LCCL_TYPE_AIV_FUNC(LCCL_ALL_REDUCE_FUNC_AUTO_DEF)"); - CheckContains(kernelTuPath, kernelTu, "LCCL_TYPE_AIV_FUNC(LCCL_REDUCE_SCATTER_FUNC_AUTO_DEF)"); - CheckContains(kernelTuPath, kernelTu, "LCCL_BROADCAST_FUNC_AUTO_DEF()"); - CheckContains(kernelTuPath, kernelTu, "LCCL_PROFILE_PROBE_FUNC_AUTO_DEF()"); - - const std::string perfTraceKernelPath = "src/collectives/kernels/perf_trace_kernel.h"; - const auto perfTraceKernel = ReadFile(perfTraceKernelPath); - CheckContains(perfTraceKernelPath, perfTraceKernel, "TILEXR_COLLECTIVES_ENABLE_PROFILING"); - CheckContains(perfTraceKernelPath, perfTraceKernel, "TileXRPerfStageBegin"); - CheckContains(perfTraceKernelPath, perfTraceKernel, "TileXRPerfStageEnd"); - CheckContains(perfTraceKernelPath, perfTraceKernel, "TileXRPerfAccumulateDuration"); - CheckContains(perfTraceKernelPath, perfTraceKernel, "TileXRPerfTraceEnabled"); - CheckContains(perfTraceKernelPath, perfTraceKernel, "#include \"comm_args.h\""); - CheckContains(perfTraceKernelPath, perfTraceKernel, "#include \"datacopy_gm2gm.h\""); - CheckContains(perfTraceKernelPath, perfTraceKernel, "TILEXR_PERF_TRACE_STATS_UB_OFFSET"); - CheckContains(perfTraceKernelPath, perfTraceKernel, "TILEXR_PERF_TRACE_STATS_OFFSET"); - CheckContains(perfTraceKernelPath, perfTraceKernel, "CpGM2UB"); - CheckContains(perfTraceKernelPath, perfTraceKernel, "CpUB2GM"); - CheckContains(perfTraceKernelPath, perfTraceKernel, "GetBlockNum()"); - CheckDoesNotContain(perfTraceKernelPath, perfTraceKernel, "header->statsOffset"); - CheckDoesNotContain(perfTraceKernelPath, perfTraceKernel, "header->maxCoreCount"); - CheckDoesNotContain(perfTraceKernelPath, perfTraceKernel, "header->stageCount"); - CheckDoesNotContain(perfTraceKernelPath, perfTraceKernel, "slot->count += 1"); - - const std::string profileProbePath = "src/collectives/kernels/kernels/lcal_profile_probe.cce"; - const auto profileProbe = ReadFile(profileProbePath); - CheckContains(profileProbePath, profileProbe, "TileXRProfileProbeKernel"); - CheckContains(profileProbePath, profileProbe, "PerfStageId::KERNEL_TOTAL"); - CheckContains(profileProbePath, profileProbe, "PerfStageId::CHUNK_TOTAL"); - CheckContains(profileProbePath, profileProbe, "PerfStageId::POST_SYNC"); - CheckContains(profileProbePath, profileProbe, "PerfStageId::LOCAL_INPUT_TO_IPC"); - CheckContains(profileProbePath, profileProbe, "PerfStageId::FLAG_POLL_WAIT"); - CheckContains(profileProbePath, profileProbe, "PerfStageId::PEER_IPC_TO_OUTPUT"); - CheckContains(profileProbePath, profileProbe, "PerfStageId::CHUNK_BARRIER"); - CheckContains(profileProbePath, profileProbe, "TileXRPerfStageBegin"); - CheckContains(profileProbePath, profileProbe, "TileXRPerfStageEnd"); - CheckContains(profileProbePath, profileProbe, "TileXRPerfAccumulateDuration"); - CheckContains(profileProbePath, profileProbe, "TILEXR_PROFILE_PROBE_COPY_UB_BYTES"); - CheckContains(profileProbePath, profileProbe, "while (copied < remain)"); - CheckContains(profileProbePath, profileProbe, "TileXR::TILEXR_PERF_TRACE_STATS_UB_OFFSET"); - CheckDoesNotContain(profileProbePath, profileProbe, "peerMems"); - CheckDoesNotContain(profileProbePath, profileProbe, "GET_IPC_MEM_ARGS"); - - const std::string perfTraceLayoutPath = "src/include/tilexr_perf_trace.h"; - const auto perfTraceLayout = ReadFile(perfTraceLayoutPath); - CheckContains(perfTraceLayoutPath, perfTraceLayout, - "Shared standalone-collectives profiling schema"); - CheckContains(perfTraceLayoutPath, perfTraceLayout, "PerfTraceCyclesToUs"); - CheckContains(perfTraceLayoutPath, perfTraceLayout, "__CCE_IS_AICORE__"); - CheckContains(perfTraceLayoutPath, perfTraceLayout, "!defined(__CCE__) || !defined(__CCE_IS_AICORE__)"); - - const std::vector kernelFiles = CollectFiles("src/collectives/kernels"); - CheckTrue(!kernelFiles.empty(), "expected collectives kernel files to be present"); - bool sawAllGatherCce = false; - bool sawAllToAllCce = false; - bool sawAllReduceCce = false; - bool sawReduceScatterCce = false; - bool sawBroadcastCce = false; - bool sawProfileProbeCce = false; - for (const auto &path : kernelFiles) { - const auto text = ReadFile(path); - CheckDoesNotContain(path, text, "tilexr_comm.h"); - CheckDoesNotContain(path, text, "reference/ascend-transformer-boost"); - CheckDoesNotContain(path, text, "LCAL_MAX_RANK_SIZE"); - if (path.find("lcal_allgather") != std::string::npos && path.find(".cce") != std::string::npos) { - sawAllGatherCce = true; - } - if (path.find("lcal_all2all_transpose.cce") != std::string::npos) { - sawAllToAllCce = true; - } - if (path.find("lcal_allreduce") != std::string::npos && path.find(".cce") != std::string::npos) { - sawAllReduceCce = true; - } - if (path.find("lcal_reduce_scatter") != std::string::npos && path.find(".cce") != std::string::npos) { - sawReduceScatterCce = true; - } - if (path.find("lcal_broadcast") != std::string::npos && path.find(".cce") != std::string::npos) { - sawBroadcastCce = true; - } - if (path.find("lcal_profile_probe.cce") != std::string::npos) { - sawProfileProbeCce = true; - } - } - CheckTrue(sawAllGatherCce, "expected copied allgather .cce sources under src/collectives/kernels"); - CheckTrue(sawAllToAllCce, "expected copied all2all .cce source under src/collectives/kernels"); - CheckTrue(sawAllReduceCce, "expected copied allreduce .cce sources under src/collectives/kernels"); - CheckTrue(sawReduceScatterCce, "expected copied reduce_scatter .cce sources under src/collectives/kernels"); - CheckTrue(sawBroadcastCce, "expected copied broadcast .cce sources under src/collectives/kernels"); - CheckTrue(sawProfileProbeCce, "expected profile probe .cce source under src/collectives/kernels"); -} - -void TestHostRegistrationLivesInCollectives() -{ - const std::string kernelPath = "src/collectives/host/collective_kernel.cpp"; - const auto kernel = ReadFile(kernelPath); - CheckContains(kernelPath, kernel, "rtDevBinaryRegister"); - CheckContains(kernelPath, kernel, "rtFunctionRegister"); - CheckContains(kernelPath, kernel, "rtKernelLaunchWithFlagV2"); - CheckContains(kernelPath, kernel, "TILEXR_TYPE2NAME"); - CheckContains(kernelPath, kernel, "TileXRCollectivesKernelBinaryData"); - CheckContains(kernelPath, kernel, "TileXRCollectivesKernelBinarySize"); - CheckContains(kernelPath, kernel, "std::mutex"); - CheckContains(kernelPath, kernel, "TileXR::TILEXR_DATA_TYPE_INT8"); - CheckContains(kernelPath, kernel, "TileXR::TILEXR_DATA_TYPE_INT16"); - CheckContains(kernelPath, kernel, "TileXR::TILEXR_DATA_TYPE_INT32"); - CheckContains(kernelPath, kernel, "TileXR::TILEXR_DATA_TYPE_INT64"); - CheckContains(kernelPath, kernel, "TileXR::TILEXR_DATA_TYPE_FP16"); - CheckContains(kernelPath, kernel, "TileXR::TILEXR_DATA_TYPE_FP32"); - CheckContains(kernelPath, kernel, "TileXR::TILEXR_DATA_TYPE_BFP16"); - CheckContains(kernelPath, kernel, "perfTrace"); - CheckContains(kernelPath, kernel, "PreparePerfTraceLaunch"); - CheckContains(kernelPath, kernel, "GetActivePerfTraceSession"); - CheckDoesNotContain(kernelPath, kernel, "g_collectiveKernelStub"); - - const auto registeredTypes = ExtractInitializer(kernelPath, kernel, "kRegisteredCollectiveTypes"); - const std::string registeredTypesPath = kernelPath + " kRegisteredCollectiveTypes[]"; - CheckContains(registeredTypesPath, registeredTypes, "TileXR::TileXRType::ALL_GATHER"); - CheckContains(registeredTypesPath, registeredTypes, "TileXR::TileXRType::ALL2ALL"); - CheckContains(registeredTypesPath, registeredTypes, "TileXR::TileXRType::ALL_REDUCE"); - CheckContains(registeredTypesPath, registeredTypes, "TileXR::TileXRType::REDUCE_SCATTER"); - CheckContains(registeredTypesPath, registeredTypes, "TileXR::TileXRType::BROADCAST"); - CheckContains(registeredTypesPath, registeredTypes, "TileXR::TileXRType::PROFILE_PROBE"); -} - -void TestPerfTraceCycleDivisorIsA5Specific() -{ - const std::string commArgsPath = "src/include/comm_args.h"; - const auto commArgs = ReadFile(commArgsPath); - CheckContains(commArgsPath, commArgs, "PERF_CYCLE_A5"); - - const std::string commPath = "src/comm/tilexr_comm.cpp"; - const auto comm = ReadFile(commPath); - CheckContains(commPath, comm, "GetChipName() == ChipName::CHIP_910A5"); - CheckContains(commPath, comm, "ExtraFlag::PERF_CYCLE_A5"); - - const std::string sessionPath = "src/collectives/host/perf_trace_session.cpp"; - const auto session = ReadFile(sessionPath); - CheckContains(sessionPath, session, "ExtraFlag::PERF_CYCLE_A5"); - CheckDoesNotContain(sessionPath, session, "TOPO_910A5) != 0 ? 1000u : 50u"); -} - -void TestDeviceKernelArgsMatchHostLaunchAbi() -{ - const std::string collectivesPath = "src/collectives/kernels/collectives.h"; - const auto collectives = ReadFile(collectivesPath); - CheckContains(collectivesPath, collectives, "GM_ADDR perfTrace"); - CheckContains(collectivesPath, collectives, "KERNELS_ARGS_CALL()"); - CheckContains(collectivesPath, collectives, - "input, output, commArgs, len, magic, op, root, cycleCount, scale, scaleCount, offset, perfTrace"); -} - -void TestBigDataAllGatherPerfStages() -{ - const std::string path = "src/collectives/kernels/kernels/lcal_allgather_big_data.cce"; - const auto text = ReadFile(path); - CheckDoesNotContain(path, text, "PerfStageId::KERNEL_TOTAL"); - CheckContains(path, text, "PerfStageId::CHUNK_TOTAL"); - CheckContains(path, text, "PerfStageId::POST_SYNC"); - CheckContains(path, text, "PerfStageId::LOCAL_INPUT_TO_IPC"); - CheckContains(path, text, "PerfStageId::FLAG_POLL_WAIT"); - CheckContains(path, text, "PerfStageId::PEER_IPC_TO_OUTPUT"); - CheckContains(path, text, "PerfStageId::CHUNK_BARRIER"); -} - -void TestTwoNpuBigDataAllGatherPerfStages() -{ - const std::string path = "src/collectives/kernels/kernels/lcal_allgather_2npu_big_data_write.cce"; - const auto text = ReadFile(path); - CheckContains(path, text, "GM_ADDR perfTrace"); - CheckDoesNotContain(path, text, "PerfStageId::KERNEL_TOTAL"); - CheckContains(path, text, "PerfStageId::CHUNK_TOTAL"); - CheckContains(path, text, "PerfStageId::POST_SYNC"); - CheckContains(path, text, "PerfStageId::LOCAL_INPUT_TO_IPC"); - CheckContains(path, text, "PerfStageId::FLAG_POLL_WAIT"); - CheckContains(path, text, "PerfStageId::PEER_IPC_TO_OUTPUT"); - CheckContains(path, text, "PerfStageId::CHUNK_BARRIER"); -} - -void TestOtherStandaloneCollectivePerfStages() -{ - const char *stageNeedles[] = { - "PerfStageId::CHUNK_TOTAL", - "PerfStageId::POST_SYNC", - "PerfStageId::LOCAL_INPUT_TO_IPC", - "PerfStageId::FLAG_POLL_WAIT", - "PerfStageId::PEER_IPC_TO_OUTPUT", - "PerfStageId::CHUNK_BARRIER", - }; - - const char *paths[] = { - "src/collectives/kernels/kernels/lcal_allreduce_big_data.cce", - "src/collectives/kernels/kernels/lcal_allreduce_2npu_big_write.cce", - "src/collectives/kernels/kernels/lcal_reduce_scatter_big_data.cce", - "src/collectives/kernels/kernels/lcal_reduce_scatter_big_data_write.cce", - "src/collectives/kernels/kernels/lcal_broadcast_big_data.cce", - "src/collectives/kernels/kernels/lcal_broadcast_write.cce", - "src/collectives/kernels/kernels/lcal_all2all_transpose.cce", - "src/collectives/kernels/91093/all2all_hierarchy.h", - "src/collectives/kernels/91093/all2all_hierarchy_small.h", - }; - for (const char *path : paths) { - const auto text = ReadFile(path); - for (const char *stageNeedle : stageNeedles) { - CheckContains(path, text, stageNeedle); - } - } -} - -void TestStandaloneCollectiveWrappersHaveCoarsePerfTrace() -{ - const std::string path = "src/collectives/kernels/lccl_op.h"; - const auto text = ReadFile(path); - CheckContains(path, text, "TileXRCoarsePerfToken"); - CheckContains(path, text, "TILEXR_COARSE_PERF_BEGIN"); - CheckContains(path, text, "TILEXR_COARSE_PERF_END"); - - const char *wrappers[] = { - "LCCL_ALLGATHER_FUNC_AUTO_DEF", - "LCCL_BROADCAST_FUNC_AUTO_DEF", - "LCCL_ALL_REDUCE_FUNC_AUTO_DEF", - "LCCL_ALL2ALL_FUNC_AUTO_DEF", - "LCCL_REDUCE_SCATTER_FUNC_AUTO_DEF", - }; - for (const char *wrapper : wrappers) { - const auto body = ExtractInitializer(path, text, wrapper); - CheckContains(std::string(path) + " " + wrapper, body, "TILEXR_COARSE_PERF_BEGIN"); - CheckContains(std::string(path) + " " + wrapper, body, "TILEXR_COARSE_PERF_END"); - } -} - -void TestCommDoesNotOwnCollectiveRuntime() -{ - const auto commFiles = CollectFiles("src/comm"); - const std::string forbidden[] = { - "src/collectives", - "TILEXR_CCE_BIN_STR", - "RegistKernel", - "LoadMTE", - "rtFunctionRegister", - "rtDevBinaryRegister", - "rtKernelLaunchWithFlagV2", - "AscendCCLKernelArgs", - "TileXRAllGather", - "TileXRAllToAll", - "TileXRAllReduce", - "TileXRReduceScatter", - "TileXRBroadcast", - }; - - for (const auto &path : commFiles) { - const auto text = ReadFile(path); - for (const auto &needle : forbidden) { - CheckDoesNotContain(path, text, needle); - } - } -} - -} // namespace - -int main() -{ - TestCollectivesOwnsCceBuild(); - TestCollectivesKernelSourcesAreScoped(); - TestHostRegistrationLivesInCollectives(); - TestPerfTraceCycleDivisorIsA5Specific(); - TestDeviceKernelArgsMatchHostLaunchAbi(); - TestBigDataAllGatherPerfStages(); - TestTwoNpuBigDataAllGatherPerfStages(); - TestOtherStandaloneCollectivePerfStages(); - TestStandaloneCollectiveWrappersHaveCoarsePerfTrace(); - TestCommDoesNotOwnCollectiveRuntime(); - return g_failures == 0 ? 0 : 1; -} diff --git a/tests/collectives/unit/test_tilexr_collectives_tools_sources.cpp b/tests/collectives/unit/test_tilexr_collectives_tools_sources.cpp deleted file mode 100644 index 032b26fb..00000000 --- a/tests/collectives/unit/test_tilexr_collectives_tools_sources.cpp +++ /dev/null @@ -1,422 +0,0 @@ -#include - -#include -#include -#include -#include -#include - -#include "../common/int32_pattern.h" - -namespace { - -int g_failures = 0; - -std::string RepoPath(const std::string &path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -bool FileExists(const std::string &path) -{ - struct stat st {}; - return stat(RepoPath(path).c_str(), &st) == 0 && S_ISREG(st.st_mode); -} - -std::string ReadFile(const std::string &path) -{ - const std::string fullPath = RepoPath(path); - std::ifstream input(fullPath.c_str()); - if (!input.is_open()) { - std::cerr << "failed to open " << fullPath << std::endl; - ++g_failures; - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -void CheckFileExists(const std::string &path) -{ - if (!FileExists(path)) { - std::cerr << "expected " << RepoPath(path) << " to exist" << std::endl; - ++g_failures; - } -} - -void CheckContains(const std::string &path, const std::string &text, const std::string &needle) -{ - if (text.find(needle) == std::string::npos) { - std::cerr << "expected " << path << " to contain: " << needle << std::endl; - ++g_failures; - } -} - -void CheckDoesNotContain(const std::string &path, const std::string &text, const std::string &needle) -{ - const auto pos = text.find(needle); - if (pos != std::string::npos) { - std::cerr << "expected " << path << " not to contain: " << needle - << " at byte " << pos << std::endl; - ++g_failures; - } -} - -void TestCorrectnessRunnerSource() -{ - const std::string path = "tests/collectives/integration/test_tilexr_collectives_correctness.cpp"; - CheckFileExists(path); - const auto text = ReadFile(path); - CheckContains(path, text, "TileXRCommInitRankLocal"); - CheckContains(path, text, "TileXRAllGather"); - CheckContains(path, text, "TileXRAllToAll"); - CheckContains(path, text, "TILEXR_DATA_TYPE_INT32"); - CheckContains(path, text, "--rank-size"); - CheckContains(path, text, "--rank"); - CheckContains(path, text, "--count"); - CheckContains(path, text, "--first-npu"); - CheckContains(path, text, "--op"); - CheckContains(path, text, "--root"); - CheckContains(path, text, "TILEXR_BROADCAST_ROOT"); - CheckContains(path, text, "options.root"); - CheckContains(path, text, "aclrtSetDevice"); - CheckContains(path, text, "aclrtCreateStream"); - CheckContains(path, text, "aclrtMalloc"); - CheckContains(path, text, "aclrtMemcpy"); - CheckContains(path, text, "aclrtSynchronizeStream"); - CheckContains(path, text, "ExpectedAllGatherValue"); - CheckContains(path, text, "ExpectedAllToAllValue"); - CheckContains(path, text, "CanUseCollisionFreeInt32Pattern"); - CheckContains(path, text, "validationCount"); - CheckContains(path, text, "CollectiveOp::REDUCESCATTER"); - CheckContains(path, text, "../common/int32_pattern.h"); - CheckDoesNotContain(path, text, "srcRank * 1000000 + index"); - CheckDoesNotContain(path, text, "dstRank * 1000 + index"); -} - -void TestPerfToolSource() -{ - const std::string path = "tests/collectives/tilexr-tests/tilexr_collective_perf.cpp"; - CheckFileExists(path); - const auto text = ReadFile(path); - CheckContains(path, text, "--op"); - CheckContains(path, text, "--min-bytes"); - CheckContains(path, text, "--max-bytes"); - CheckContains(path, text, "--step-factor"); - CheckContains(path, text, "--iters"); - CheckContains(path, text, "--warmup-iters"); - CheckContains(path, text, "--datatype"); - CheckContains(path, text, "--rank-size"); - CheckContains(path, text, "--rank"); - CheckContains(path, text, "--first-npu"); - CheckContains(path, text, "--check"); - CheckContains(path, text, "--csv"); - CheckContains(path, text, "--min-algbw"); - CheckContains(path, text, "--max-latency-us"); - CheckContains(path, text, "--profile"); - CheckContains(path, text, "--profile-dir"); - CheckContains(path, text, "--profile-ai-prompt"); - CheckContains(path, text, "--profile-sample-every"); - CheckContains(path, text, "--comm-mode"); - CheckContains(path, text, "TileXRCommInitRank("); - CheckContains(path, text, "TileXRGetUniqueId("); - CheckContains(path, text, "WriteProfileHostInfo"); - CheckContains(path, text, "TILEXR_PROFILE_CLOCK_OFFSET_NS"); - CheckContains(path, text, "clock_offset_ns"); - CheckContains(path, text, "clock_reference"); - CheckContains(path, text, "epoch_ns"); - CheckContains(path, text, "--op allgather|alltoall|allreduce|reducescatter|broadcast|profile-probe"); - CheckContains(path, text, "CollectiveOp::NOOP"); - CheckContains(path, text, "CollectiveOp::PROFILE_PROBE"); - CheckContains(path, text, "TileXRProfileProbe"); - CheckContains(path, text, "TileXRAllReduce"); - CheckContains(path, text, "TileXRReduceScatter"); - CheckContains(path, text, "TileXRBroadcast"); - CheckContains(path, text, "reducescatter: count * rank_size * dtype_size"); - CheckContains(path, text, "broadcast: count * dtype_size"); - CheckContains(path, text, "aclrtCreateEvent"); - CheckContains(path, text, "aclrtRecordEvent"); - CheckContains(path, text, "aclrtEventElapsedTime"); - CheckContains(path, text, "algbw(GB/s)"); - CheckContains(path, text, "busbw(GB/s)"); - CheckContains(path, text, "avg(us)"); - CheckContains(path, text, "min(us)"); - CheckContains(path, text, "max(us)"); - CheckContains(path, text, "ParseDataType"); - CheckContains(path, text, "int8"); - CheckContains(path, text, "int16"); - CheckContains(path, text, "int32"); - CheckContains(path, text, "int64"); - CheckContains(path, text, "fp16"); - CheckContains(path, text, "fp32"); - CheckContains(path, text, "bf16"); - CheckContains(path, text, "ComputeAlgBandwidthGbps"); - CheckContains(path, text, "ComputeBusBandwidthGbps"); - CheckContains(path, text, "ValidateInt32"); - CheckContains(path, text, "CheckedMulInt64"); - CheckContains(path, text, "CheckedBytesForElements"); - CheckContains(path, text, "kMaxHostBufferBytes"); - CheckContains(path, text, "FitsHostBufferLimit"); - CheckContains(path, text, "AdvanceBytes"); - CheckContains(path, text, "aclrtMemcpy H2D devRecv sentinel"); - CheckContains(path, text, "actualSendBytesPerRank"); - CheckContains(path, text, "CanUseCollisionFreeInt32Pattern"); - CheckContains(path, text, "../common/int32_pattern.h"); - CheckContains(path, text, "JoinPath"); - CheckContains(path, text, "ResolveProfileOutputDir"); - CheckContains(path, text, "\"run/prof/collectives\""); - CheckContains(path, text, "\"rank\" + std::to_string(options.rank)"); - CheckContains(path, text, "\"launch\" + std::to_string(profileLaunchIndex)"); - CheckContains(path, text, "TileXRCollectivePerfSessionCreate"); - CheckContains(path, text, "TileXRCollectivePerfSetActiveSession"); - CheckContains(path, text, "TileXRCollectivePerfWriteReport"); - CheckContains(path, text, "TileXRCollectivePerfSessionDestroy"); - CheckContains(path, text, "aclrtSynchronizeStream before perf report"); - CheckContains(path, text, "ProfileThisLaunch"); - CheckContains(path, text, "profileLaunchIndex"); - CheckContains(path, text, "TileXRCollectivePerfSetActiveSession(perfSession)"); - CheckContains(path, text, "TileXRCollectivePerfSetActiveSession(nullptr)"); - CheckContains(path, text, "skipWriteReport"); - CheckContains(path, text, "TileXRCollectivePerfSetActiveSession clear failed"); - CheckContains(path, text, "TileXRCollectivePerfSessionDestroy failed"); - CheckDoesNotContain(path, text, "sendElements * static_cast(options.dtype.bytes)"); - CheckDoesNotContain(path, text, "static_cast(static_cast(bytes) * options.stepFactor)"); - CheckDoesNotContain(path, text, "srcRank * 1000000 + index"); - CheckDoesNotContain(path, text, "dstRank * 1000 + index"); -} - -void TestInt32PatternHasNoKnownCollisions() -{ - using TileXRCollectivesTest::CanUseCollisionFreeInt32Pattern; - using TileXRCollectivesTest::ExpectedAllGatherValue; - using TileXRCollectivesTest::ExpectedAllToAllValue; - - const int rankSize = 8; - const int64_t reviewedCollisionA = 9646; - const int64_t reviewedCollisionB = 76425; - if (ExpectedAllGatherValue(rankSize, 0, reviewedCollisionA) == - ExpectedAllGatherValue(rankSize, 0, reviewedCollisionB)) { - std::cerr << "reviewed allgather collision was reintroduced" << std::endl; - ++g_failures; - } - - if (!CanUseCollisionFreeInt32Pattern(128, 262144)) { - std::cerr << "pattern should cover 1MiB INT32 validation at rank_size=128" << std::endl; - ++g_failures; - } - - std::set seen; - const int64_t sampleCount = 4096; - for (int src = 0; src < rankSize; ++src) { - for (int64_t i = 0; i < sampleCount; ++i) { - const int32_t value = ExpectedAllGatherValue(rankSize, src, i); - if (!seen.insert(value).second) { - std::cerr << "allgather int32 pattern collision src=" << src << " index=" << i << std::endl; - ++g_failures; - return; - } - } - } - - seen.clear(); - for (int src = 0; src < rankSize; ++src) { - for (int dst = 0; dst < rankSize; ++dst) { - for (int64_t i = 0; i < sampleCount; ++i) { - const int32_t value = ExpectedAllToAllValue(rankSize, src, dst, i); - if (!seen.insert(value).second) { - std::cerr << "alltoall int32 pattern collision src=" << src - << " dst=" << dst << " index=" << i << std::endl; - ++g_failures; - return; - } - } - } - } -} - -void TestLauncherScripts() -{ - const std::string correctnessPath = "tests/collectives/run_collectives_correctness.sh"; - CheckFileExists(correctnessPath); - const auto correctness = ReadFile(correctnessPath); - CheckContains(correctnessPath, correctness, "rank_size"); - CheckContains(correctnessPath, correctness, "count"); - CheckContains(correctnessPath, correctness, "first_npu"); - CheckContains(correctnessPath, correctness, "bin_dir"); - CheckContains(correctnessPath, correctness, "collectives_correctness_rank${rank}.log"); - CheckContains(correctnessPath, correctness, "TILEXR_SKIP_IF_INSUFFICIENT_NPUS"); - CheckContains(correctnessPath, correctness, "TILEXR_COLLECTIVES_RUN_TIMEOUT_SEC"); - CheckContains(correctnessPath, correctness, "kill_remaining_children"); - CheckContains(correctnessPath, correctness, "tail_logs"); - CheckContains(correctnessPath, correctness, "Timed out after"); - CheckContains(correctnessPath, correctness, "watchdog_pid"); - CheckContains(correctnessPath, correctness, "sleep \"${timeout_sec}\" >/dev/null 2>&1 &"); - CheckContains(correctnessPath, correctness, "wait \"${watchdog_pid}\""); - CheckContains(correctnessPath, correctness, "wait -n"); - CheckContains(correctnessPath, correctness, "completed_count"); - CheckDoesNotContain(correctnessPath, correctness, "wait -n -p completed_pid"); - CheckContains(correctnessPath, correctness, "npu-smi info -l"); - CheckContains(correctnessPath, correctness, "tail -n"); - CheckContains(correctnessPath, correctness, "test_tilexr_collectives_correctness"); - - const std::string perfPath = "tools/collectives/run_collective_perf.sh"; - CheckFileExists(perfPath); - const auto perf = ReadFile(perfPath); - CheckContains(perfPath, perf, "tilexr_collective_perf"); - CheckContains(perfPath, perf, "--rank-size"); - CheckContains(perfPath, perf, "--rank"); - CheckContains(perfPath, perf, "--first-npu"); - CheckContains(perfPath, perf, "collective_perf_rank${rank}.log"); - CheckContains(perfPath, perf, "TILEXR_SKIP_IF_INSUFFICIENT_NPUS"); - CheckContains(perfPath, perf, "TILEXR_COLLECTIVES_RUN_TIMEOUT_SEC"); - CheckContains(perfPath, perf, "kill_remaining_children"); - CheckContains(perfPath, perf, "tail_logs"); - CheckContains(perfPath, perf, "Timed out after"); - CheckContains(perfPath, perf, "watchdog_pid"); - CheckContains(perfPath, perf, "sleep \"${timeout_sec}\" >/dev/null 2>&1 &"); - CheckContains(perfPath, perf, "wait \"${watchdog_pid}\""); - CheckContains(perfPath, perf, "wait -n"); - CheckContains(perfPath, perf, "completed_count"); - CheckDoesNotContain(perfPath, perf, "wait -n -p completed_pid"); - CheckContains(perfPath, perf, "parse_profile_args"); - CheckContains(perfPath, perf, "write_profile_report_if_enabled"); - CheckContains(perfPath, perf, "tilexr_collective_profile_report.py"); - CheckContains(perfPath, perf, "tools/collectives"); - CheckContains(perfPath, perf, "--warmup-iters"); - CheckContains(perfPath, perf, "--profile-sample-every"); - CheckContains(perfPath, perf, "is_true_bool"); - CheckContains(perfPath, perf, "yes"); - - CheckFileExists("tools/collectives/tilexr_collective_profile_report.py"); - CheckFileExists("tools/collectives/run_collective_perf.sh"); - CheckFileExists("tools/collectives/run_collective_perf_multihost.sh"); - CheckFileExists("tools/collectives/run_collective_perf_multihost_suite.sh"); - - const std::string multiHostPerfPath = "tools/collectives/run_collective_perf_multihost.sh"; - CheckFileExists(multiHostPerfPath); - const auto multiHostPerf = ReadFile(multiHostPerfPath); - CheckContains(multiHostPerfPath, multiHostPerf, "TILEXR_MULTIHOST_PEERS"); - CheckContains(multiHostPerfPath, multiHostPerf, "TILEXR_COMM_ID"); - CheckContains(multiHostPerfPath, multiHostPerf, "--comm-mode socket"); - CheckContains(multiHostPerfPath, multiHostPerf, "tilexr_collective_profile_report.py"); - CheckContains(multiHostPerfPath, multiHostPerf, "tools/collectives"); - CheckContains(multiHostPerfPath, multiHostPerf, "copy_rank_profile"); - CheckContains(multiHostPerfPath, multiHostPerf, "test -d '${profile_dir}/rank${rank}'"); - CheckContains(multiHostPerfPath, multiHostPerf, "remote_epoch_ns"); - CheckContains(multiHostPerfPath, multiHostPerf, "TILEXR_MULTIHOST_REMOTE_REPO_DIR"); - CheckContains(multiHostPerfPath, multiHostPerf, "rank0_host_ip"); - CheckContains(multiHostPerfPath, multiHostPerf, "comm_id=\"${TILEXR_COMM_ID:-${rank0_host_ip}:10067}\""); - CheckContains(multiHostPerfPath, multiHostPerf, "TILEXR_PROFILE_CLOCK_OFFSET_NS"); - CheckContains(multiHostPerfPath, multiHostPerf, "TILEXR_PROFILE_CLOCK_SYNC_REFERENCE"); - CheckContains(multiHostPerfPath, multiHostPerf, "command -v rsync"); - CheckContains(multiHostPerfPath, multiHostPerf, "tar -cf - ."); - CheckContains(multiHostPerfPath, multiHostPerf, "tar -xf - -C"); - CheckContains(multiHostPerfPath, multiHostPerf, "rank${rank}"); - CheckContains(multiHostPerfPath, multiHostPerf, "host_label=\"${target#*@}\""); - CheckContains(multiHostPerfPath, multiHostPerf, "cd \"${repo_dir}\""); - CheckContains(multiHostPerfPath, multiHostPerf, "build_dir=\"$(cd \"${bin_dir}/../..\" && pwd)\""); - CheckContains(multiHostPerfPath, multiHostPerf, "${build_dir}/src/collectives"); - CheckContains(multiHostPerfPath, multiHostPerf, "clock_offset_ns=0"); - CheckContains(multiHostPerfPath, multiHostPerf, "reference_midpoint_ns"); - CheckContains(multiHostPerfPath, multiHostPerf, "set +u"); - CheckContains(multiHostPerfPath, multiHostPerf, "ASCEND_PROCESS_LOG_PATH"); - CheckContains(multiHostPerfPath, multiHostPerf, "plog/rank${rank}"); -} - -void TestCMakeWiring() -{ - const std::string path = "tests/collectives/CMakeLists.txt"; - const auto text = ReadFile(path); - CheckContains(path, text, "integration/test_tilexr_collectives_correctness.cpp"); - CheckContains(path, text, "tilexr-tests/tilexr_collective_perf.cpp"); - CheckContains(path, text, "add_executable(test_tilexr_collectives_correctness"); - CheckContains(path, text, "add_executable(tilexr_collective_perf"); - CheckContains(path, text, "target_link_libraries(test_tilexr_collectives_correctness"); - CheckContains(path, text, "target_link_libraries(tilexr_collective_perf"); - CheckContains(path, text, "TileXRAddAscendRpathLink"); - CheckContains(path, text, "LINKER:-rpath-link,${ASCEND_HOME_PATH}/${ARCH}-linux/lib64"); - CheckContains(path, text, "LINKER:-rpath-link,${ASCEND_DRIVER_PATH}/lib64/driver"); - CheckContains(path, text, "test_tilexr_collectives_tools_sources"); - CheckContains(path, text, "run_collectives_correctness.sh"); - CheckContains(path, text, "../../tools/collectives/run_collective_perf.sh"); - CheckContains(path, text, "../../tools/collectives/run_collective_perf_multihost.sh"); - CheckContains(path, text, "../../tools/collectives/run_collective_perf_multihost_suite.sh"); - CheckContains(path, text, "find_package(Python3 COMPONENTS Interpreter)"); - CheckContains(path, text, "test_collective_profile_report"); - CheckContains(path, text, "tilexr_collective_profile_report.py"); - CheckContains(path, text, "../../tools/collectives/tilexr_collective_profile_report.py"); - CheckDoesNotContain(path, text, "add_test(NAME test_tilexr_collectives_correctness"); - CheckDoesNotContain(path, text, "add_test(NAME tilexr_collective_perf"); -} - -void TestReadmeDocumentsManualRuns() -{ - const std::string path = "tests/collectives/README.md"; - CheckFileExists(path); - const auto text = ReadFile(path); - CheckContains(path, text, "libtilexr-collectives"); - CheckContains(path, text, "tile-comm"); - CheckContains(path, text, "test_tilexr_collectives_correctness"); - CheckContains(path, text, "tilexr_collective_perf"); - CheckContains(path, text, "run_collectives_correctness.sh"); - CheckContains(path, text, "run_collective_perf.sh"); - CheckContains(path, text, "--rank-size"); - CheckContains(path, text, "--rank"); - CheckContains(path, text, "--first-npu"); - CheckContains(path, text, "--op"); - CheckContains(path, text, "--datatype"); - CheckContains(path, text, "--min-bytes"); - CheckContains(path, text, "--max-bytes"); - CheckContains(path, text, "--check"); - CheckContains(path, text, "--csv"); - CheckContains(path, text, "algbw(GB/s)"); - CheckContains(path, text, "busbw(GB/s)"); - CheckContains(path, text, "actual send bytes per rank"); - CheckContains(path, text, "--op allgather|alltoall|allreduce|reducescatter|broadcast|profile-probe"); - CheckContains(path, text, - "Message-size semantics: allgather/allreduce/broadcast: count * dtype_size; alltoall/reducescatter: count * rank_size * dtype_size."); - CheckContains(path, text, "Checked"); - CheckContains(path, text, "allreduce/reducescatter runs require `--datatype int32`"); - CheckContains(path, text, "rank_size - 1"); - CheckContains(path, text, "trace_index.json"); - CheckContains(path, text, "perfetto_trace.json"); - CheckContains(path, text, "root-level report.html"); - CheckContains(path, text, "rank-level summary"); - CheckContains(path, text, "launch/rank/stage event names"); - CheckContains(path, text, "Multi-Host Profiling"); - CheckContains(path, text, "../../docs/COLLECTIVES_MULTIHOST_PROFILING.md"); - CheckContains(path, text, "TILEXR_MULTIHOST_PEERS"); - CheckContains(path, text, "--comm-mode socket"); - CheckContains(path, text, "--op profile-probe"); - CheckContains(path, text, "tilexr_collective_profile_report.py"); - CheckContains(path, text, "../../tools/collectives/tilexr_collective_profile_report.py"); - CheckContains(path, text, "zoomable chronological timeline"); - CheckContains(path, text, "warmup launches are not profiled"); - CheckContains(path, text, "TILEXR_COLLECTIVES_RUN_TIMEOUT_SEC"); - CheckContains(path, text, "TILEXR_SKIP_IF_INSUFFICIENT_NPUS"); - CheckContains(path, text, "CTest"); - CheckContains(path, text, "manual"); -} - -} // namespace - -int main() -{ - TestCorrectnessRunnerSource(); - TestPerfToolSource(); - TestLauncherScripts(); - TestCMakeWiring(); - TestReadmeDocumentsManualRuns(); - TestInt32PatternHasNoKnownCollisions(); - if (g_failures != 0) { - std::cerr << g_failures << " collectives tools source checks failed" << std::endl; - return 1; - } - std::cout << "TileXR collectives tools source checks passed" << std::endl; - return 0; -} diff --git a/tests/collectives/unit/test_vllm_collectives_integration_sources.py b/tests/collectives/unit/test_vllm_collectives_integration_sources.py deleted file mode 100644 index 7b7ba4ee..00000000 --- a/tests/collectives/unit/test_vllm_collectives_integration_sources.py +++ /dev/null @@ -1,374 +0,0 @@ -#!/usr/bin/env python3 -import subprocess -from pathlib import Path - - -ROOT = Path(__file__).resolve().parents[3] - - -def read_rel(path: str) -> str: - return (ROOT / path).read_text(encoding="utf-8") - - -def assert_exists(path: str) -> None: - full_path = ROOT / path - assert full_path.exists(), f"missing required integration file: {path}" - - -def test_vllm_ascend_shim_files_exist() -> None: - for path in [ - "integrations/vllm_ascend/tilexr_collectives/__init__.py", - "integrations/vllm_ascend/tilexr_collectives/runtime.py", - "integrations/vllm_ascend/tilexr_collectives/torch_collectives.py", - "integrations/vllm_ascend/tilexr_collectives/vllm_patch.py", - "integrations/vllm_ascend/smoke_collectives.py", - "integrations/vllm_ascend/run_tilexr_collectives_smoke.sh", - "integrations/vllm_ascend/sitecustomize.py", - "tests/collectives/deploy_and_run_vllm_remote.sh", - ]: - assert_exists(path) - - -def test_runtime_uses_tilexr_c_abi_and_not_hccl() -> None: - runtime_source = read_rel("integrations/vllm_ascend/tilexr_collectives/runtime.py") - required_tokens = [ - "TileXRCommInitRankLocal", - "TileXRCommDestroy", - "TileXRAllGather", - "TileXRAllToAll", - "TILEXR_DATA_TYPE_INT32", - "TILEXR_DATA_TYPE_FP16", - ] - for token in required_tokens: - assert token in runtime_source - assert "hccl" not in runtime_source.lower() - - -def test_runtime_exposes_phase3_collective_c_api() -> None: - runtime_source = read_rel("integrations/vllm_ascend/tilexr_collectives/runtime.py") - for token in [ - "TileXRAllReduce", - "TileXRReduceScatter", - "TileXRBroadcast", - "TILEXR_REDUCE_SUM", - "def all_reduce(", - "def reduce_scatter(", - "def broadcast(", - ]: - assert token in runtime_source - - -def test_torch_helpers_require_contiguous_npu_tensors() -> None: - source = read_rel("integrations/vllm_ascend/tilexr_collectives/torch_collectives.py") - for token in [ - "def all_gather", - "def all_to_all", - "tensor.is_contiguous()", - "tensor.device.type != \"npu\"", - "torch.npu.current_stream()", - "torch.npu.set_device(device_index)", - ]: - assert token in source - assert "torch.uint8" not in source - assert "TILEXR_DATA_TYPE_UINT8" not in source - - -def test_torch_helpers_expose_vllm_compatible_collectives() -> None: - source = read_rel("integrations/vllm_ascend/tilexr_collectives/torch_collectives.py") - for token in [ - "def _move_dim_to_front(", - "def _restore_dim_from_front(", - "def all_gather(tensor, rank: int, world_size: int, install_prefix: str, dim: int = -1", - "def all_reduce(tensor, rank: int, world_size: int, install_prefix: str", - "def reduce_scatter(tensor, rank: int, world_size: int, install_prefix: str, dim: int = -1", - "def broadcast(tensor, rank: int, world_size: int, install_prefix: str, root: int = 0", - ]: - assert token in source - - -def test_vllm_adapter_is_opt_in_and_import_lightweight() -> None: - source = read_rel("integrations/vllm_ascend/tilexr_collectives/vllm_adapter.py") - for token in [ - "VLLM_ASCEND_TILEXR_COLLECTIVES", - "class TileXRVllmCollectivesAdapter", - "def enabled(", - "def all_reduce(", - "def all_gather(", - "def reduce_scatter(", - "def all_to_all(", - "def broadcast(", - "root=src", - "def should_fallback(", - ]: - assert token in source - assert "import vllm" not in source - assert "from vllm" not in source - - -def test_vllm_adapter_normalizes_all_to_all_dims_before_fallback() -> None: - source = read_rel("integrations/vllm_ascend/tilexr_collectives/vllm_adapter.py") - assert "_same_normalized_dim(" in source - assert "if not self._same_normalized_dim(input_, scatter_dim, gather_dim):" in source - assert "if scatter_dim != gather_dim:" not in source - - -def test_vllm_patch_hooks_npu_communicator_methods() -> None: - source = read_rel("integrations/vllm_ascend/tilexr_collectives/vllm_patch.py") - for token in [ - "def patch_npu_communicator(", - "VLLM_ASCEND_TILEXR_COLLECTIVES", - "TILEXR_VLLM_TRACE", - "NPUCommunicator", - "_tilexr_collectives_adapter", - "_tilexr_collectives_route_counts", - "all_reduce", - "all_gather", - "reduce_scatter", - "all_to_all", - "broadcast", - "return original(self", - ]: - assert token in source - - -def test_sitecustomize_auto_patches_only_when_enabled() -> None: - source = read_rel("integrations/vllm_ascend/sitecustomize.py") - for token in [ - "VLLM_ASCEND_TILEXR_COLLECTIVES", - "patch_npu_communicator", - "sitecustomize", - ]: - assert token in source - assert "raise" not in source - - -def test_remote_script_is_isolated_and_logs_environment() -> None: - source = read_rel("tests/collectives/deploy_and_run_vllm_remote.sh") - for token in [ - "TILEXR_VLLM_REMOTE", - "TILEXR_VLLM_REMOTE_BASE", - "TILEXR_VLLM_REMOTE_ASCEND_HOME_PATH", - "TILEXR_VLLM_REMOTE_ASCEND_DRIVER_PATH", - "TILEXR_VLLM_REMOTE_CMAKE_CCE_COMPILER", - "rsync -a --delete", - "sync_local_submodule", - "--exclude='.worktrees'", - "npu-smi info", - "run_tilexr_collectives_smoke.sh", - ]: - assert token in source - forbidden_tokens = [ - "TILEXR_VLLM_REMOTE:-", - "TILEXR_VLLM_REMOTE_BASE:-", - "REMOTE_BASE=/", - "REMOTE_BASE=\"/", - "submodule update --init --recursive", - ">> ~/.bashrc", - "pip install --user", - "apt-get install", - "yum install", - "rm -rf /usr/local/Ascend", - ] - for token in forbidden_tokens: - assert token not in source - assert "3rdparty/hcomm" not in source - - -def test_remote_script_supports_selected_python_environment() -> None: - source = read_rel("tests/collectives/deploy_and_run_vllm_remote.sh") - for token in [ - "TILEXR_VLLM_REMOTE_PYTHON", - "TILEXR_VLLM_REMOTE_CONDA_ENV", - "TILEXR_VLLM_REMOTE_CONDA_SH", - "TILEXR_VLLM_REMOTE_PYTHONPATH", - "select_remote_python", - "selected_python=", - "remote_pythonpath", - "dump_selected_python_environment", - "run_selected_python_preflight", - "torch.npu.current_stream()", - "npu_stream", - "TILEXR_VLLM_SMOKE_PYTHON=\"\\${selected_python}\"", - "TILEXR_VLLM_SMOKE_PYTHONPATH=\"\\${remote_pythonpath}\"", - ]: - assert token in source - assert "python3 -m pip show torch" not in source - assert "python3 -m pip show torch-npu" not in source - - -def test_remote_script_can_pass_consistent_ssh_options_to_ssh_and_rsync() -> None: - source = read_rel("tests/collectives/deploy_and_run_vllm_remote.sh") - for token in [ - "TILEXR_VLLM_REMOTE_SSH_OPTS", - "ssh_args=(", - "rsync_ssh_cmd=(", - "ssh \"${ssh_args[@]}\"", - "rsync -e \"${rsync_ssh_cmd[*]}\"", - ]: - assert token in source - - -def test_remote_script_can_probe_vllm_source_trees_without_crashing() -> None: - source = read_rel("tests/collectives/deploy_and_run_vllm_remote.sh") - for token in [ - "TILEXR_VLLM_REMOTE_VLLM_SOURCE", - "TILEXR_VLLM_REMOTE_VLLM_ASCEND_SOURCE", - "TILEXR_VLLM_REMOTE_DUMMY_MODEL", - "TILEXR_VLLM_REMOTE_VLLM_PLUGINS", - "TILEXR_VLLM_REMOTE_NNAL_ATB_SET_ENV", - "build_vllm_probe_pythonpath", - "run_vllm_import_probe", - "subprocess.run", - "TORCH_DEVICE_BACKEND_AUTOLOAD", - "vllm.distributed.device_communicators.base_device_communicator", - "vllm_ascend.distributed.device_communicators.npu_communicator", - "probe_vllm_environment \"pre-cann\"", - "probe_vllm_environment \"post-cann\"", - "local probe_label=\"\\${1:?probe label required}\"", - "pythonpath_entries+=(\"\\${remote_pythonpath}\")", - "probe_vllm_communicator_patch", - "TileXR vllm NPUCommunicator patch probe", - "TileXR vllm NPUCommunicator route counts", - "patch_npu_communicator(\"install\")", - "probe_vllm_dummy_inference", - "load_format=\"dummy\"", - "skip_tokenizer_init=True", - "PASS TileXR vllm dummy inference probe", - "source \"\\${remote_nnal_atb_set_env}\"", - ]: - assert token in source - assert "local probe_label=\"${1:?probe label required}\"" not in source - - -def test_smoke_launcher_supports_python_override() -> None: - source = read_rel("integrations/vllm_ascend/run_tilexr_collectives_smoke.sh") - for token in [ - "PYTHON_BIN=\"${TILEXR_VLLM_SMOKE_PYTHON:-python3}\"", - "command -v \"${PYTHON_BIN}\"", - "import sys", - "ERROR: Python command failed interpreter preflight", - "TileXR vllm collectives smoke", - "\"${PYTHON_BIN}\" \"${SCRIPT_DIR}/smoke_collectives.py\"", - ]: - assert token in source - assert "python3 \"${SCRIPT_DIR}/smoke_collectives.py\"" not in source - - -def test_smoke_launcher_accepts_preflight_sentinel_with_plugin_stdout(tmp_path) -> None: - fake_python = tmp_path / "fake-python" - fake_python.write_text( - """#!/usr/bin/env bash -set -euo pipefail -if [[ "${1:-}" == "-c" ]]; then - if [[ "${2:-}" == *"tilexr_python_preflight_ok"* ]]; then - echo "INFO plugin stdout before sentinel" - echo "tilexr_python_preflight_ok" - exit 0 - fi - if [[ "${2:-}" == *"sys.executable"* ]]; then - printf '%s\\n' "$0" - exit 0 - fi -fi -if [[ "${1:-}" == *"smoke_collectives.py" ]]; then - exit 0 -fi -exit 0 -""", - encoding="utf-8", - ) - fake_python.chmod(0o755) - - launcher = ROOT / "integrations/vllm_ascend/run_tilexr_collectives_smoke.sh" - result = subprocess.run( - [ - "bash", - str(launcher), - "2", - "16", - "0", - str(tmp_path / "install"), - "allgather", - "int32", - ], - env={ - "PATH": "/usr/bin:/bin", - "TILEXR_VLLM_SMOKE_PYTHON": str(fake_python), - "TILEXR_VLLM_SMOKE_TIMEOUT_SEC": "5", - }, - text=True, - stdout=subprocess.PIPE, - stderr=subprocess.PIPE, - timeout=10, - ) - - assert result.returncode == 0, result.stdout + result.stderr - assert "PASS TileXR vllm collectives smoke" in result.stdout - - -def test_phase3_smoke_covers_core_collectives() -> None: - smoke_source = read_rel("integrations/vllm_ascend/smoke_collectives.py") - launcher_source = read_rel("integrations/vllm_ascend/run_tilexr_collectives_smoke.sh") - remote_source = read_rel("tests/collectives/deploy_and_run_vllm_remote.sh") - for token in ["allreduce", "reducescatter", "broadcast"]: - assert token in smoke_source - assert token in launcher_source - assert token in remote_source - for token in ["probe_vllm_environment", "vllm_ascend", "VLLM_ASCEND_TILEXR_COLLECTIVES"]: - assert token in remote_source - - -def test_phase3_docs_describe_feature_flag_and_boundaries() -> None: - readme = read_rel("tests/collectives/README.md") - for token in [ - "VLLM_ASCEND_TILEXR_COLLECTIVES=1", - "TILEXR_VLLM_REMOTE_VLLM_SOURCE", - "TILEXR_VLLM_REMOTE_VLLM_ASCEND_SOURCE", - "TILEXR_VLLM_REMOTE_DUMMY_MODEL", - "TILEXR_VLLM_REMOTE_NNAL_ATB_SET_ENV", - "TILEXR_VLLM_REMOTE_PYTHONPATH", - "tilexr-vllm29", - "torch==2.9.0", - "torch-npu==2.9.0", - "triton-ascend==3.2.0", - "vllm_ascend.distributed.device_communicators.npu_communicator", - "NPUCommunicator patch probe", - "TileXR vllm NPUCommunicator route counts", - "TileXR vllm dummy inference probe", - "NPUCommunicator patched", - "load_format=\"dummy\"", - "libatb.so", - "torchvision", - "rc=0", - "allreduce", - "reducescatter", - "broadcast", - "fallback", - "libhccl.so", - "zmq", - "vllm-ascend inference", - ]: - assert token in readme - - -def main() -> None: - test_vllm_ascend_shim_files_exist() - test_runtime_uses_tilexr_c_abi_and_not_hccl() - test_runtime_exposes_phase3_collective_c_api() - test_torch_helpers_require_contiguous_npu_tensors() - test_torch_helpers_expose_vllm_compatible_collectives() - test_vllm_adapter_is_opt_in_and_import_lightweight() - test_vllm_adapter_normalizes_all_to_all_dims_before_fallback() - test_vllm_patch_hooks_npu_communicator_methods() - test_sitecustomize_auto_patches_only_when_enabled() - test_remote_script_is_isolated_and_logs_environment() - test_remote_script_supports_selected_python_environment() - test_remote_script_can_probe_vllm_source_trees_without_crashing() - test_smoke_launcher_supports_python_override() - test_phase3_smoke_covers_core_collectives() - test_phase3_docs_describe_feature_flag_and_boundaries() - print("PASS vllm collectives integration source checks") - - -if __name__ == "__main__": - main() diff --git a/tests/comm/CMakeLists.txt b/tests/comm/CMakeLists.txt index 76a7a553..ab8d6c8e 100644 --- a/tests/comm/CMakeLists.txt +++ b/tests/comm/CMakeLists.txt @@ -31,17 +31,8 @@ target_include_directories(test_tilexr_log_spdlog_compile PRIVATE ${TILEXR_ROOT}/src/comm ) -add_executable(test_tilexr_source_guards - unit/test_tilexr_source_guards.cpp -) - -target_compile_definitions(test_tilexr_source_guards PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - install(TARGETS test_tilexr_log test_tilexr_log_spdlog_compile - test_tilexr_source_guards RUNTIME DESTINATION ${CMAKE_INSTALL_PREFIX}/bin ) diff --git a/tests/comm/unit/test_tilexr_source_guards.cpp b/tests/comm/unit/test_tilexr_source_guards.cpp deleted file mode 100644 index 4c7e70e0..00000000 --- a/tests/comm/unit/test_tilexr_source_guards.cpp +++ /dev/null @@ -1,216 +0,0 @@ -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -std::string RepoPath(const std::string& path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -std::string ReadFile(const std::string& path) -{ - std::ifstream input(RepoPath(path).c_str()); - if (!input.is_open()) { - std::cerr << "failed to open " << RepoPath(path) << std::endl; - ++g_failures; - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -void CheckContains(const std::string& path, const std::string& text, const std::string& needle) -{ - if (text.find(needle) == std::string::npos) { - std::cerr << "expected text not found in " << path << ": " << needle << std::endl; - ++g_failures; - } -} - -void CheckNotContains(const std::string& path, const std::string& text, const std::string& needle) -{ - const auto pos = text.find(needle); - if (pos != std::string::npos) { - std::cerr << "unexpected text in " << path << ": " << needle - << " at byte " << pos << std::endl; - ++g_failures; - } -} - -std::string RunCommand(const std::string& command) -{ - FILE* pipe = popen(command.c_str(), "r"); - if (pipe == nullptr) { - std::cerr << "failed to run command: " << command << std::endl; - ++g_failures; - return {}; - } - - std::string output; - char buffer[4096]; - while (fgets(buffer, sizeof(buffer), pipe) != nullptr) { - output += buffer; - } - - const int status = pclose(pipe); - if (status != 0) { - std::cerr << "command failed: " << command << std::endl; - ++g_failures; - } - return output; -} - -void CheckNoLineContains(const std::string& description, const std::string& text, const std::string& needle) -{ - std::istringstream lines(text); - std::string line; - while (std::getline(lines, line)) { - if (line.find(needle) != std::string::npos) { - std::cerr << "unexpected " << description << ": " << line << std::endl; - ++g_failures; - } - } -} - -void TestOpenSourceTarballsAreNotTracked() -{ - const std::string command = "git -C " + RepoPath(".") + " ls-files 3rdparty/open_source"; - const auto trackedFiles = RunCommand(command); - - CheckNoLineContains("tracked open-source dependency archive", trackedFiles, ".tar.gz"); - CheckNoLineContains("tracked open-source dependency archive", trackedFiles, ".tar.xz"); -} - -void TestCommInitChecksDeviceCommArgsSync() -{ - const std::string path = "src/comm/tilexr_comm.cpp"; - const auto text = ReadFile(path); - - CheckNotContains(path, text, "InitMem();\n g_localPeerMemMap"); - CheckNotContains(path, text, " SyncCommArgs();"); - CheckContains(path, text, "ret = SyncCommArgs();"); - CheckContains(path, text, "ret = InitMem();"); -} - -void TestCWrappersDoNotPublishFailedCommunicators() -{ - const std::string path = "src/comm/comm_wrap.cpp"; - const auto text = ReadFile(path); - - CheckNotContains(path, text, "*comm = c;\n int ret = c->Init();"); - CheckNotContains(path, text, "*comm = c;\n int ret = c->InitThread"); - CheckNotContains(path, text, "comms[i] = new (std::nothrow) TileXRComm"); - CheckContains(path, text, "c.release()"); - CheckContains(path, text, "commHolders[i].release()"); -} - -void TestDumpInitCleansFailedAllocations() -{ - const std::string path = "src/comm/tilexr_comm.cpp"; - const auto text = ReadFile(path); - - CheckContains(path, text, "aclrtFree(dumpAddr);"); - CheckContains(path, text, "std::free(memory);"); -} - -void TestSocketExchangeUsesDirectConnectionsOnly() -{ - const std::string cppPath = "src/comm/tools/socket/tilexr_sock_exchange.cpp"; - const std::string headerPath = "src/comm/tools/socket/tilexr_sock_exchange.h"; - const auto cppText = ReadFile(cppPath); - const auto headerText = ReadFile(headerPath); - - CheckNotContains(cppPath, cppText, "StartSecureTunnel"); - CheckNotContains(headerPath, headerText, "StartSecureTunnel"); - CheckNotContains(cppPath, cppText, "popen"); - CheckNotContains(cppPath, cppText, "/usr/bin/ssh"); - CheckNotContains(headerPath, headerText, "FILE* pipe_"); - CheckNotContains(headerPath, headerText, "lockFileDescriptor_"); - CheckContains(cppPath, cppText, "return Connect();"); - CheckContains(cppPath, cppText, "bool envProvided = false;"); - CheckContains(cppPath, cppText, "if (!envProvided) {"); - CheckNotContains(cppPath, cppText, - "handle->addr.sin.sin_port = htons(TILEXR_DEFAULT_SOCK_PORT + dev + commDomain);"); -} - -void TestRuntimeEnvDoesNotPrependCannDevlib() -{ - const std::string path = "scripts/common_env.sh"; - const auto text = ReadFile(path); - - CheckNotContains(path, text, "${ASCEND_HOME_PATH}/${TILEXR_OS_ARCH}-linux/devlib"); -} - -void TestRuntimeEnvUsesReadableDriverShim() -{ - const std::string path = "scripts/common_env.sh"; - const auto text = ReadFile(path); - - CheckContains(path, text, "TILEXR_DRIVER_SHIM_HOME"); - CheckContains(path, text, "[ ! -r \"${ASCEND_DRIVER_PATH}/kernel/inc\" ]"); - CheckContains(path, text, "${ASCEND_HOME_PATH}/${TILEXR_OS_ARCH}-linux/include/driver"); - CheckContains(path, text, "export ASCEND_DRIVER_PATH=${TILEXR_DRIVER_SHIM_HOME}"); -} - -void TestRootCMakeRespectsAscendDriverOverride() -{ - const std::string path = "CMakeLists.txt"; - const auto text = ReadFile(path); - - CheckNotContains(path, text, "set(ASCEND_DRIVER_PATH /usr/local/Ascend/driver)"); - CheckContains(path, text, "set(_tilexr_default_ascend_driver_path \"$ENV{ASCEND_DRIVER_PATH}\")"); - CheckContains(path, text, "set(ASCEND_DRIVER_PATH \"${_tilexr_default_ascend_driver_path}\" CACHE PATH"); -} - -void TestCommBuildIncludesProfilingHeaders() -{ - const std::string rootPath = "CMakeLists.txt"; - const std::string commPath = "src/comm/CMakeLists.txt"; - const auto rootText = ReadFile(rootPath); - const auto commText = ReadFile(commPath); - - CheckContains(rootPath, rootText, "${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/profiling/"); - CheckContains(commPath, commText, "${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/profiling/"); -} - -void TestChipMapCoversObservedAscend950Variants() -{ - const std::string path = "src/comm/tilexr_internal.cpp"; - const auto text = ReadFile(path); - - CheckContains(path, text, "{\"Ascend950PR\", ChipName::CHIP_950PR}"); - CheckContains(path, text, "{\"Ascend950PR_9589\", ChipName::CHIP_950PR}"); - CheckContains(path, text, "{\"Ascend950PR_9599\", ChipName::CHIP_950PR}"); -} - -} // namespace - -int main() -{ - TestOpenSourceTarballsAreNotTracked(); - TestCommInitChecksDeviceCommArgsSync(); - TestCWrappersDoNotPublishFailedCommunicators(); - TestDumpInitCleansFailedAllocations(); - TestSocketExchangeUsesDirectConnectionsOnly(); - TestRuntimeEnvDoesNotPrependCannDevlib(); - TestRootCMakeRespectsAscendDriverOverride(); - TestCommBuildIncludesProfilingHeaders(); - TestChipMapCoversObservedAscend950Variants(); - - if (g_failures != 0) { - std::cerr << g_failures << " TileXR source guard checks failed" << std::endl; - return 1; - } - std::cout << "TileXR source guard checks passed" << std::endl; - return 0; -} diff --git a/tests/data_as_flag/CMakeLists.txt b/tests/data_as_flag/CMakeLists.txt index d3138aed..0ffd573a 100644 --- a/tests/data_as_flag/CMakeLists.txt +++ b/tests/data_as_flag/CMakeLists.txt @@ -21,19 +21,9 @@ target_include_directories(test_tilexr_data_as_flag_header_compile PRIVATE ${TILEXR_ROOT}/src/include ) -add_executable(test_tilexr_data_as_flag_source_guard - unit/test_tilexr_data_as_flag_source_guard.cpp -) - -target_compile_definitions(test_tilexr_data_as_flag_source_guard PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - add_test(NAME test_tilexr_data_as_flag_header_compile COMMAND test_tilexr_data_as_flag_header_compile) -add_test(NAME test_tilexr_data_as_flag_source_guard COMMAND test_tilexr_data_as_flag_source_guard) install(TARGETS test_tilexr_data_as_flag_header_compile - test_tilexr_data_as_flag_source_guard RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} ) diff --git a/tests/data_as_flag/unit/test_tilexr_data_as_flag_source_guard.cpp b/tests/data_as_flag/unit/test_tilexr_data_as_flag_source_guard.cpp deleted file mode 100644 index d8038a0c..00000000 --- a/tests/data_as_flag/unit/test_tilexr_data_as_flag_source_guard.cpp +++ /dev/null @@ -1,135 +0,0 @@ -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -std::string RepoPath(const std::string& path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -std::string ReadFile(const std::string& path) -{ - const std::string fullPath = RepoPath(path); - std::ifstream input(fullPath.c_str()); - if (!input.is_open()) { - std::cerr << "failed to open " << fullPath << std::endl; - ++g_failures; - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -void CheckContains(const std::string& path, const std::string& text, const std::string& needle) -{ - if (text.find(needle) == std::string::npos) { - std::cerr << "expected " << path << " to contain: " << needle << std::endl; - ++g_failures; - } -} - -void CheckDoesNotContain(const std::string& path, const std::string& text, const std::string& needle) -{ - const auto pos = text.find(needle); - if (pos != std::string::npos) { - std::cerr << "expected " << path << " not to contain: " << needle - << " at byte " << pos << std::endl; - ++g_failures; - } -} - -void TestHeaderShape() -{ - const std::string path = "src/include/tilexr_data_as_flag.h"; - const auto text = ReadFile(path); - CheckContains(path, text, "DATA_AS_FLAG_BLOCK_BYTES = 512"); - CheckContains(path, text, "DATA_AS_FLAG_PAYLOAD_BYTES = 480"); - CheckContains(path, text, "DATA_AS_FLAG_FLAG_BYTES = 32"); - CheckContains(path, text, "DATA_AS_FLAG_FLAG_OFFSET_BYTES = 480"); - CheckContains(path, text, "DATA_AS_FLAG_FLAG_FLOATS"); - CheckContains(path, text, "DATA_AS_FLAG_READY_VALUE = 1.0f"); - CheckContains(path, text, "DataAsFlagBlockCountForPayloadBytes"); - CheckContains(path, text, "DataAsFlagInit"); - CheckContains(path, text, "DataAsFlagSend"); - CheckContains(path, text, "DataAsFlagCheck"); - CheckContains(path, text, "DataAsFlagCheckAndRecv"); -} - -void TestHeaderUsesExpectedAscendCApis() -{ - const std::string path = "src/include/tilexr_data_as_flag.h"; - const auto text = ReadFile(path); - CheckContains(path, text, "#include \"adv_api/reduce/sum.h\""); - CheckContains(path, text, "AscendC::DataCopyPad"); - CheckContains(path, text, "AscendC::Sum<"); - CheckContains(path, text, "AscendC::SumParams"); - CheckContains(path, text, "sharedTmpBuffer"); - CheckContains(path, text, "DataAsFlagSumWorkspaceBytes"); - CheckContains(path, text, "ReinterpretCast"); - CheckContains(path, text, "GetSize()"); - CheckContains(path, text, "DATA_AS_FLAG_SUM_RESULT_BYTES"); - CheckContains(path, text, "DataAsFlagMaxRecvBlocks"); - CheckContains(path, text, "while (!DataAsFlagCheckBatch"); - CheckDoesNotContain(path, text, "GlobalTensor::GetValue"); - CheckDoesNotContain(path, text, "GlobalTensor::SetValue"); - CheckDoesNotContain(path, text, "checkScratchBlockCapacity"); -} - -void TestInstallWiring() -{ - const std::string path = "src/comm/CMakeLists.txt"; - const auto text = ReadFile(path); - CheckContains(path, text, "${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_data_as_flag.h"); - CheckContains(path, text, "DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}"); -} - -void TestRootCMakeWiring() -{ - const std::string path = "CMakeLists.txt"; - const auto text = ReadFile(path); - CheckContains(path, text, "add_subdirectory(tests/data_as_flag)"); -} - -void TestDataAsFlagCMakeWiring() -{ - const std::string path = "tests/data_as_flag/CMakeLists.txt"; - const auto text = ReadFile(path); - CheckContains(path, text, "add_executable(test_tilexr_data_as_flag_header_compile"); - CheckContains(path, text, "unit/test_tilexr_data_as_flag_header_compile.cpp"); - CheckContains(path, text, "add_executable(test_tilexr_data_as_flag_source_guard"); - CheckContains(path, text, "unit/test_tilexr_data_as_flag_source_guard.cpp"); - CheckContains(path, text, "target_include_directories(test_tilexr_data_as_flag_header_compile PRIVATE"); - CheckContains(path, text, "target_compile_definitions(test_tilexr_data_as_flag_source_guard PRIVATE"); - CheckContains(path, text, "add_test(NAME test_tilexr_data_as_flag_header_compile"); - CheckContains(path, text, "add_test(NAME test_tilexr_data_as_flag_source_guard"); - CheckContains(path, text, "install(TARGETS"); - CheckContains(path, text, "test_tilexr_data_as_flag_header_compile"); - CheckContains(path, text, "test_tilexr_data_as_flag_source_guard"); -} - -} // namespace - -int main() -{ - TestHeaderShape(); - TestHeaderUsesExpectedAscendCApis(); - TestInstallWiring(); - TestRootCMakeWiring(); - TestDataAsFlagCMakeWiring(); - if (g_failures != 0) { - std::cerr << g_failures << " DataAsFlag source guard checks failed" << std::endl; - return 1; - } - std::cout << "TileXR DataAsFlag source guard checks passed" << std::endl; - return 0; -} diff --git a/tests/ep/CMakeLists.txt b/tests/ep/CMakeLists.txt index e505f1d8..02ac7d59 100644 --- a/tests/ep/CMakeLists.txt +++ b/tests/ep/CMakeLists.txt @@ -79,12 +79,6 @@ add_executable(test_tilexr_ep_layout ${TILEXR_ROOT}/src/ep/host/ep_layout.cpp ) -add_executable(test_tilexr_ep_api_sources - unit/test_tilexr_ep_api_sources.cpp -) - -add_executable(test_tilexr_ep_kernel_sources unit/test_tilexr_ep_kernel_sources.cpp) - add_executable(test_tilexr_ep_host_validation unit/test_tilexr_ep_host_validation.cpp ${TILEXR_ROOT}/src/ep/host/ep_layout.cpp @@ -99,18 +93,11 @@ target_include_directories(test_tilexr_ep_host_validation PRIVATE ${TILEXR_EP_TEST_INCLUDE_DIRS} ) -target_compile_definitions(test_tilexr_ep_api_sources PRIVATE TILEXR_SOURCE_ROOT="${TILEXR_ROOT}") -target_compile_definitions(test_tilexr_ep_kernel_sources PRIVATE TILEXR_SOURCE_ROOT="${TILEXR_ROOT}") - add_test(NAME test_tilexr_ep_layout COMMAND test_tilexr_ep_layout) -add_test(NAME test_tilexr_ep_api_sources COMMAND test_tilexr_ep_api_sources) -add_test(NAME test_tilexr_ep_kernel_sources COMMAND test_tilexr_ep_kernel_sources) add_test(NAME test_tilexr_ep_host_validation COMMAND test_tilexr_ep_host_validation) install(TARGETS test_tilexr_ep_layout - test_tilexr_ep_api_sources - test_tilexr_ep_kernel_sources test_tilexr_ep_host_validation RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} ) diff --git a/tests/ep/README.md b/tests/ep/README.md index cace06ef..081010dc 100644 --- a/tests/ep/README.md +++ b/tests/ep/README.md @@ -2,20 +2,18 @@ This tree tests the standalone TileXR EP module under `src/ep`. It is independent from hcomm, HCCL window helpers, and `ops-transformer`; the same-node route uses TileXR IPC peer-memory windows and `SyncCollectives`, while cross-node dispatch/combine use TileXR-registered UDMA workspaces. -## Source-Only Tests +## Host Tests From `tests/ep`: ```bash source ../../scripts/common_env.sh -bash build.sh source-only +bash build.sh host-only ./install/bin/test_tilexr_ep_layout -./install/bin/test_tilexr_ep_api_sources ./install/bin/test_tilexr_ep_host_validation -./install/bin/test_tilexr_ep_kernel_sources ``` -`source-only` mode builds and installs the source-layout, API source, host validation, and kernel source tests without building the hardware demo. +`host-only` mode builds and installs the layout and host validation tests without building the hardware demo. ## Full Hardware Demo diff --git a/tests/ep/build.sh b/tests/ep/build.sh index a4732f34..4fb12e8f 100755 --- a/tests/ep/build.sh +++ b/tests/ep/build.sh @@ -3,7 +3,7 @@ set -euo pipefail SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" TILEXR_ROOT="$(cd "${SCRIPT_DIR}/../.." && pwd)" -MODE="${1:-source-only}" +MODE="${1:-host-only}" : "${ASCEND_HOME_PATH:=}" : "${LD_LIBRARY_PATH:=}" @@ -19,11 +19,11 @@ case "${MODE}" in cmake --build "${ROOT_BUILD_DIR}" --target install -j"$(nproc)" DEMO_OPTION="-DBUILD_TILEXR_EP_DEMO=ON" ;; - source-only) + host-only) DEMO_OPTION="-DBUILD_TILEXR_EP_DEMO=OFF" ;; *) - echo "Usage: $0 [source-only|full]" >&2 + echo "Usage: $0 [host-only|full]" >&2 exit 2 ;; esac diff --git a/tests/ep/unit/test_tilexr_ep_api_sources.cpp b/tests/ep/unit/test_tilexr_ep_api_sources.cpp deleted file mode 100644 index 7f942804..00000000 --- a/tests/ep/unit/test_tilexr_ep_api_sources.cpp +++ /dev/null @@ -1,306 +0,0 @@ -#include -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -#ifdef TILEXR_SOURCE_ROOT -const char *kSourceRoot = TILEXR_SOURCE_ROOT; -#else -const char *kSourceRoot = "."; -#endif - -std::string JoinPath(const std::string &base, const std::string &path) -{ - if (base.empty() || base[base.size() - 1] == '/') { - return base + path; - } - return base + "/" + path; -} - -bool ReadFile(const std::string &relativePath, std::string *contents) -{ - const std::string fullPath = JoinPath(kSourceRoot, relativePath); - std::ifstream stream(fullPath.c_str()); - if (!stream.is_open()) { - std::cerr << "missing file: " << relativePath << std::endl; - ++g_failures; - return false; - } - - std::ostringstream buffer; - buffer << stream.rdbuf(); - *contents = buffer.str(); - return true; -} - -void CheckContains(const std::string &label, const std::string &contents, const std::string &needle) -{ - if (contents.find(needle) == std::string::npos) { - std::cerr << label << " missing: " << needle << std::endl; - ++g_failures; - } -} - -void CheckNotContains(const std::string &label, const std::string &contents, const std::string &needle) -{ - if (contents.find(needle) != std::string::npos) { - std::cerr << label << " contains forbidden string: " << needle << std::endl; - ++g_failures; - } -} - -void TestPublicHeader() -{ - std::string contents; - if (!ReadFile("src/include/tilexr_ep.h", &contents)) { - return; - } - - CheckContains("src/include/tilexr_ep.h", contents, "#ifdef __cplusplus"); - CheckContains("src/include/tilexr_ep.h", contents, "extern \"C\""); - CheckContains("src/include/tilexr_ep.h", contents, "int TileXRMoeEpDispatch("); - CheckContains("src/include/tilexr_ep.h", contents, "int TileXRMoeEpDispatchV2("); - CheckContains("src/include/tilexr_ep.h", contents, "int TileXRMoeEpCombine("); - CheckContains("src/include/tilexr_ep.h", contents, "int TileXRMoeEpCombineV2("); - CheckContains("src/include/tilexr_ep.h", contents, "TileXRCommPtr comm"); - CheckContains("src/include/tilexr_ep.h", contents, "TileXR::TileXRDataType dtype"); - CheckContains("src/include/tilexr_ep.h", contents, "aclrtStream stream"); - CheckContains("src/include/tilexr_ep.h", contents, "xActiveMask"); - CheckContains("src/include/tilexr_ep.h", contents, "dynamicScalesOut"); - CheckContains("src/include/tilexr_ep.h", contents, "tpRecvCountsOut"); - CheckContains("src/include/tilexr_ep.h", contents, "expandScalesOut"); - CheckContains("src/include/tilexr_ep.h", contents, "workspace"); - CheckContains("src/include/tilexr_ep.h", contents, "quantMode"); - CheckContains("src/include/tilexr_ep.h", contents, "tpWorldSize"); - CheckContains("src/include/tilexr_ep.h", contents, "sharedExpertNum"); - - std::string apiContents; - if (ReadFile("src/include/tilexr_api.h", &apiContents)) { - CheckContains("src/include/tilexr_api.h", apiContents, "TileXRGetUDMARegistryHost"); - } -} - -void TestBuildPlacement() -{ - std::string rootCmake; - if (ReadFile("CMakeLists.txt", &rootCmake)) { - CheckContains("CMakeLists.txt", rootCmake, - "option(TILEXR_BUILD_EP \"Build TileXR EP communication library\" OFF)"); - CheckContains("CMakeLists.txt", rootCmake, "add_subdirectory(src/ep)"); - } - - std::string epCmake; - if (ReadFile("src/ep/CMakeLists.txt", &epCmake)) { - CheckContains("src/ep/CMakeLists.txt", epCmake, "add_library(tilexr-ep SHARED"); - CheckContains("src/ep/CMakeLists.txt", epCmake, "tile-comm"); - CheckContains("src/ep/CMakeLists.txt", epCmake, "libtilexr_ep_combine_kernel.so"); - CheckContains("src/ep/CMakeLists.txt", epCmake, "tilexr_ep.h"); - CheckContains("src/ep/CMakeLists.txt", epCmake, "install(TARGETS tilexr-ep"); - } -} - -void TestEpHostChecksRegisteredWorkspace() -{ - std::string launchContext; - if (!ReadFile("src/ep/host/ep_launch_context.cpp", &launchContext)) { - return; - } - - CheckContains("src/ep/host/ep_launch_context.cpp", launchContext, "ValidateRegisteredWorkspace"); - CheckContains("src/ep/host/ep_launch_context.cpp", launchContext, "TileXRGetUDMARegistryHost"); - CheckContains("src/ep/host/ep_launch_context.cpp", launchContext, "UDMARegionContains"); - CheckContains("src/ep/host/ep_launch_context.cpp", launchContext, "TileXREpUdmaRequiredWorkspaceBytes"); -} - -void TestEpSocDefaultFollowsEnvironment() -{ - std::string epCmake; - if (!ReadFile("src/ep/CMakeLists.txt", &epCmake)) { - return; - } - - CheckContains("src/ep/CMakeLists.txt", epCmake, "$ENV{TILEXR_SOC_NAME}"); - CheckContains("src/ep/CMakeLists.txt", epCmake, "string(TOLOWER"); - CheckContains("src/ep/CMakeLists.txt", epCmake, "ascend910b"); - CheckContains("src/ep/CMakeLists.txt", epCmake, "dav-c220-vec"); - CheckContains("src/ep/CMakeLists.txt", epCmake, "Ascend910B"); -} - -void TestChipMapRecognizesAscend950Dt9582() -{ - std::string internal; - if (!ReadFile("src/comm/tilexr_internal.cpp", &internal)) { - return; - } - - CheckContains("src/comm/tilexr_internal.cpp", internal, "\"Ascend950DT_9582\", ChipName::CHIP_950"); - CheckContains("src/comm/tilexr_internal.cpp", internal, "\"Ascend950PR_9599\", ChipName::CHIP_950PR"); - CheckContains("src/comm/tilexr_internal.cpp", internal, "bool UseLegacyIpcPid(ChipName chipName)"); - CheckContains("src/comm/tilexr_internal.cpp", internal, "chipName == ChipName::CHIP_950PR"); -} - -void TestEpKernelUsesCceArchFlags() -{ - std::string epCmake; - if (!ReadFile("src/ep/CMakeLists.txt", &epCmake)) { - return; - } - - CheckContains("src/ep/CMakeLists.txt", epCmake, "-xcce"); - CheckContains("src/ep/CMakeLists.txt", epCmake, "${TILEXR_EP_AICORE_ARCH}"); - CheckContains("src/ep/CMakeLists.txt", epCmake, "set(TILEXR_EP_KERNEL_LINK_OPTIONS ${TILEXR_EP_AICORE_ARCH})"); - CheckNotContains("src/ep/CMakeLists.txt", epCmake, "-xasc"); - CheckNotContains("src/ep/CMakeLists.txt", epCmake, "--npu-arch="); - CheckNotContains("src/ep/CMakeLists.txt", epCmake, "--cce-auto-infer-kernel-type=false"); - CheckNotContains("src/ep/CMakeLists.txt", epCmake, "--cce-fatobj-link"); -} - -void TestSyncFlagBuffersUseFullFlagUnit() -{ - std::string syncHeader; - if (!ReadFile("src/include/tilexr_sync.h", &syncHeader)) { - return; - } - - CheckContains("src/include/tilexr_sync.h", syncHeader, - "LocalTensor localSet = tBuf.GetWithOffset(FLAG_UNIT_INT_NUM, 0);"); - CheckContains("src/include/tilexr_sync.h", syncHeader, - "LocalTensor localWait = tBuf.GetWithOffset(FLAG_UNIT_INT_NUM, 0);"); -} - -const char *kRemoteDeployScript = "tests/ep/demo/deploy_and_run_remote.sh"; - -void TestRemoteDeployScriptCleansRemoteCheckout() -{ - std::string deployScript; - if (!ReadFile(kRemoteDeployScript, &deployScript)) { - return; - } - - CheckContains(kRemoteDeployScript, deployScript, "case \"\\${remote_repo}\" in"); - CheckContains(kRemoteDeployScript, deployScript, "Refusing to clean unexpected remote repo"); - CheckContains(kRemoteDeployScript, deployScript, "rm -rf -- \"\\${remote_repo}\""); - CheckContains(kRemoteDeployScript, deployScript, "mkdir -p -- \"\\${remote_repo}\""); -} - -void TestRemoteDeployScriptDoesNotInitializeSubmodules() -{ - std::string deployScript; - if (!ReadFile(kRemoteDeployScript, &deployScript)) { - return; - } - - CheckNotContains(kRemoteDeployScript, deployScript, "submodule update"); - CheckNotContains(kRemoteDeployScript, deployScript, "3rdparty/hcomm"); - CheckNotContains(kRemoteDeployScript, deployScript, "ops-transformer"); - CheckNotContains(kRemoteDeployScript, deployScript, "3rdparty/shmem"); - CheckNotContains(kRemoteDeployScript, deployScript, "3rdparty/spdlog"); -} - -void TestRemoteDeployScriptDoesNotExposePrivateRemoteDefaults() -{ - std::string deployScript; - if (!ReadFile(kRemoteDeployScript, &deployScript)) { - return; - } - - CheckContains(kRemoteDeployScript, deployScript, "TILEXR_EP_REMOTE:?set TILEXR_EP_REMOTE"); - CheckContains(kRemoteDeployScript, deployScript, "TILEXR_EP_REMOTE_BASE:?set TILEXR_EP_REMOTE_BASE"); - CheckNotContains(kRemoteDeployScript, deployScript, "TILEXR_EP_REMOTE:-"); - CheckNotContains(kRemoteDeployScript, deployScript, "TILEXR_EP_REMOTE_BASE:-"); - CheckNotContains(kRemoteDeployScript, deployScript, "REMOTE_BASE=/"); -} - -void TestDemoRunnerUsesLibAndLib64Paths() -{ - std::string runner; - if (!ReadFile("tests/ep/demo/run_tilexr_ep_dispatch_demo.sh", &runner)) { - return; - } - - CheckContains("tests/ep/demo/run_tilexr_ep_dispatch_demo.sh", runner, "${TILEXR_ROOT}/install/lib64"); - CheckContains("tests/ep/demo/run_tilexr_ep_dispatch_demo.sh", runner, "${TILEXR_ROOT}/install/lib"); - CheckContains("tests/ep/demo/run_tilexr_ep_dispatch_demo.sh", runner, "${INSTALL_DIR}/lib64"); - CheckContains("tests/ep/demo/run_tilexr_ep_dispatch_demo.sh", runner, "${INSTALL_DIR}/lib"); -} - -void TestDispatchDemoRegistersAlignedUdmaWorkspace() -{ - std::string demo; - if (!ReadFile("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", &demo)) { - return; - } - - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "AlignAddress"); - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "rawWorkspaceDev"); - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, - "workspaceDev = reinterpret_cast(AlignAddress("); - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, - "TileXRUDMARegister(comm, static_cast(workspaceDev), workspaceBytes"); - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "EpRequiredWorkspaceBytes"); -} - -void TestDispatchDemoUsesHostBarrierBeforeValidation() -{ - std::string demo; - if (!ReadFile("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", &demo)) { - return; - } - - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "DemoBarrierAll"); - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "TILEXR_DEMO_BARRIER_ADDR"); - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "dispatch synchronized"); -} - -void TestNoForbiddenDependencies() -{ - const std::vector paths = { - "src/include/tilexr_ep.h", - "src/ep/CMakeLists.txt", - "src/ep/host/ep_layout.h", - "src/ep/host/ep_layout.cpp", - }; - const std::vector forbidden = { - "ops-transformer", - "GetHcclContext", - "TileXRUDMARegister", - "UDMAPut", - "shmem", - }; - - for (std::vector::const_iterator path = paths.begin(); path != paths.end(); ++path) { - std::string contents; - if (!ReadFile(*path, &contents)) { - continue; - } - for (std::vector::const_iterator needle = forbidden.begin(); needle != forbidden.end(); ++needle) { - CheckNotContains(*path, contents, *needle); - } - } -} - -} // namespace - -int main() -{ - TestPublicHeader(); - TestBuildPlacement(); - TestEpHostChecksRegisteredWorkspace(); - TestEpSocDefaultFollowsEnvironment(); - TestChipMapRecognizesAscend950Dt9582(); - TestEpKernelUsesCceArchFlags(); - TestSyncFlagBuffersUseFullFlagUnit(); - TestRemoteDeployScriptCleansRemoteCheckout(); - TestRemoteDeployScriptDoesNotInitializeSubmodules(); - TestRemoteDeployScriptDoesNotExposePrivateRemoteDefaults(); - TestDemoRunnerUsesLibAndLib64Paths(); - TestDispatchDemoRegistersAlignedUdmaWorkspace(); - TestDispatchDemoUsesHostBarrierBeforeValidation(); - TestNoForbiddenDependencies(); - return g_failures == 0 ? 0 : 1; -} diff --git a/tests/ep/unit/test_tilexr_ep_kernel_sources.cpp b/tests/ep/unit/test_tilexr_ep_kernel_sources.cpp deleted file mode 100644 index a8100561..00000000 --- a/tests/ep/unit/test_tilexr_ep_kernel_sources.cpp +++ /dev/null @@ -1,478 +0,0 @@ -#include -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -#ifdef TILEXR_SOURCE_ROOT -const char *kSourceRoot = TILEXR_SOURCE_ROOT; -#else -const char *kSourceRoot = "."; -#endif - -std::string JoinPath(const std::string &base, const std::string &path) -{ - if (base.empty() || base[base.size() - 1] == '/') { - return base + path; - } - return base + "/" + path; -} - -bool ReadFile(const std::string &relativePath, std::string *contents) -{ - const std::string fullPath = JoinPath(kSourceRoot, relativePath); - std::ifstream stream(fullPath.c_str()); - if (!stream.is_open()) { - std::cerr << "missing file: " << relativePath << std::endl; - ++g_failures; - return false; - } - - std::ostringstream buffer; - buffer << stream.rdbuf(); - *contents = buffer.str(); - return true; -} - -void CheckContains(const std::string &label, const std::string &contents, const std::string &needle) -{ - if (contents.find(needle) == std::string::npos) { - std::cerr << label << " missing: " << needle << std::endl; - ++g_failures; - } -} - -void CheckNotContains(const std::string &label, const std::string &contents, const std::string &needle) -{ - if (contents.find(needle) != std::string::npos) { - std::cerr << label << " contains forbidden string: " << needle << std::endl; - ++g_failures; - } -} - -void TestKernelUsesTileXRPeerMemory() -{ - const std::string path = "src/ep/kernels/tilexr_ep_dispatch_kernel.cpp"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckContains(path, contents, "extern \"C\" __global__ __aicore__ void tilexr_ep_dispatch_kernel"); - CheckContains(path, contents, "launch_tilexr_ep_dispatch_kernel"); - CheckContains(path, contents, "CommArgs"); - CheckContains(path, contents, "peerMems"); - CheckContains(path, contents, "GlobalTensor peerMems"); - CheckContains(path, contents, "peerMems.GetValue(peer)"); - CheckContains(path, contents, "IPC_DATA_OFFSET"); - CheckContains(path, contents, "SyncCollectives"); - CheckContains(path, contents, "DataCopyPad"); - CheckContains(path, contents, "kEpStepWindowCleared"); - CheckContains(path, contents, "kEpStepDispatchReady"); - CheckContains(path, contents, "kEpStepDispatchDrained"); - CheckContains(path, contents, "LoadInt32FromGm"); - CheckContains(path, contents, "LoadAssistTupleFromGm"); - CheckContains(path, contents, "StoreWindowHeader"); - CheckContains(path, contents, "StoreSlotHeader"); - CheckContains(path, contents, "StoreAssistTuple"); - CheckContains(path, contents, "TileXREpFlushDispatchSlotHeaders"); - CheckContains(path, contents, "sourceWindow = TileXREpWindowBase"); - CheckNotContains(path, contents, "expertIds["); - CheckNotContains(path, contents, "assistBase[item]"); - CheckNotContains(path, contents, "args->peerMems[peer]"); - CheckNotContains(path, contents, "slot->count"); - CheckNotContains(path, contents, "assist[index]"); -} - -void TestCrossNodeDispatchUsesUDMARegistry() -{ - const std::string path = "src/ep/kernels/tilexr_ep_dispatch_kernel.cpp"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckContains(path, contents, "tilexr_udma.h"); - CheckContains(path, contents, "TileXR::UDMARegistryEnabled(args)"); - CheckContains(path, contents, "TileXREpUsesUdmaWindow"); - CheckContains(path, contents, "TileXR::UDMAPutNbi"); - CheckContains(path, contents, "TileXR::UDMAQuiet(args, dstRank)"); - CheckContains(path, contents, "TileXREpFlushDispatchSlotHeaders"); - CheckContains(path, contents, "sourceWindow = TileXREpWindowBase"); -} - -void TestCrossNodeDispatchPullsRemoteSlots() -{ - const std::string path = "src/ep/kernels/tilexr_ep_dispatch_kernel.cpp"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckContains(path, contents, "tilexr_ep_dispatch_cross_node_kernel"); - CheckContains(path, contents, "launch_tilexr_ep_dispatch_cross_node_kernel"); - CheckContains(path, contents, "TileXREpPullUdmaSlots"); - CheckContains(path, contents, "TileXR::UDMAGetNbi"); - CheckContains(path, contents, "TileXREpNotifyUdmaReady"); - CheckContains(path, contents, "TileXREpWaitUdmaReady"); - CheckContains(path, contents, "TileXREpNotifyAllUdmaReady"); - CheckContains(path, contents, "TileXREpWaitAllUdmaReady"); - CheckContains(path, contents, "TileXR::UDMAPutSignalNbi"); -} - -void TestCrossNodeDispatchSeparatesLocalAndRemotePeers() -{ - const std::string path = "src/ep/kernels/tilexr_ep_dispatch_kernel.cpp"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckContains(path, contents, "TileXREpIsSameNodePeer"); - CheckContains(path, contents, "TileXREpUsesUdmaPeer"); - CheckContains(path, contents, "if (localRankSize <= 1)"); - CheckContains(path, contents, "TileXREpPublishLocalUdmaSlot"); - CheckContains(path, contents, "TileXREpDispatchWriteWindow"); - CheckContains(path, contents, "args->localRankSize"); - CheckContains(path, contents, "if (localRankSize > 1)"); - CheckContains(path, contents, "dstRank != rank && TileXREpIsSameNodePeer(rank, dstRank, localRankSize)"); - CheckContains(path, contents, "srcRank != rank &&"); - CheckContains(path, contents, "TileXREpIsSameNodePeer(rank, srcRank, localRankSize)"); - CheckContains(path, contents, "sameNodeSource ? rank : srcRank"); - CheckContains(path, contents, "!TileXREpIsSameNodePeer(rank, peer, localRankSize)"); -} - -void TestHostDispatchSplitsCrossNodeKernel() -{ - const std::string path = "src/ep/host/ep_kernel_launch.cpp"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckContains(path, contents, "launch_tilexr_ep_dispatch_cross_node_kernel"); - CheckContains(path, contents, "TileXREpUsesCrossNodeKernel"); -} - -void TestDispatchHelpersLiveInDispatchHelperFile() -{ - const std::string path = "src/ep/kernels/tilexr_ep_dispatch_helpers.h"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckContains(path, contents, "TileXREpRouteToDstRank"); - CheckContains(path, contents, "TileXREpCopyRoutePayload"); - CheckContains(path, contents, "TileXREpStoreDispatchSlotHeader"); - CheckContains(path, contents, "TileXREpStoreAssistTuple"); - CheckContains(path, contents, "localWindow + PayloadOffset(dstRank, slotBytes)"); -} - -void TestCombineHelpersLiveInCombineHelperFile() -{ - const std::string path = "src/ep/kernels/tilexr_ep_combine_helpers.h"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckContains(path, contents, "TileXREpDrainSourceWindow"); - CheckContains(path, contents, "TileXREpWaitDispatchSlotReady"); - CheckContains(path, contents, "slotMagic"); - CheckContains(path, contents, "TileXREpLoadAssistTuple"); - CheckContains(path, contents, "TileXREpGetCombineTokenId"); - CheckContains(path, contents, "TileXREpGetCombineTopKId"); - CheckContains(path, contents, "sourceWindow + SlotOffset(slotRank, slotBytes)"); -} - -void TestCombineKernelUsesTileXRPeerMemory() -{ - const std::string path = "src/ep/kernels/tilexr_ep_combine_kernel.cpp"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckContains(path, contents, "extern \"C\" __global__ __aicore__ void tilexr_ep_combine_kernel"); - CheckContains(path, contents, "launch_tilexr_ep_combine_kernel"); - CheckContains(path, contents, "kEpStepCombineWindowCleared"); - CheckContains(path, contents, "kEpStepCombineReady"); - CheckContains(path, contents, "ScatterCombineRows"); - CheckContains(path, contents, "DrainCombineRows"); - CheckContains(path, contents, "AccumulateRow"); - CheckContains(path, contents, "tilexr_ep_combine_cross_node_kernel"); - CheckContains(path, contents, "tilexr_ep_combine_cross_node_drain_kernel"); - CheckContains(path, contents, "TileXREpNotifyRemoteUdmaReadySeparate"); - CheckContains(path, contents, "TileXREpWaitRemoteUdmaReady"); - CheckNotContains(path, contents, "tilexr_ep_dispatch_kernel"); - - std::string hostLaunch; - if (ReadFile("src/ep/host/ep_kernel_launch.cpp", &hostLaunch)) { - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "TileXREpLaunchCombineKernel"); - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "launch_tilexr_ep_combine_kernel"); - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, - "launch_tilexr_ep_combine_cross_node_kernel"); - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "aclrtSynchronizeStream"); - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "TILEXR_ERROR_TIMEOUT"); - } -} - -void TestKernelCommonHasCombineHelpers() -{ - const std::string path = "src/ep/kernels/tilexr_ep_kernel_common.h"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckContains(path, contents, "UDMASecondOperationOffset"); - CheckContains(path, contents, "TileXREpNotifyRemoteUdmaReadySeparate"); - CheckContains(path, contents, "TileXREpWaitRemoteUdmaReady"); - CheckContains(path, contents, "TileXREpStoreStatusValue"); - CheckContains(path, contents, "TileXREpFlushUdmaSourceWindow"); - CheckContains(path, contents, "IsValidShape"); - - std::string combine; - if (ReadFile("src/ep/kernels/tilexr_ep_combine_kernel.cpp", &combine)) { - CheckContains("src/ep/kernels/tilexr_ep_combine_kernel.cpp", combine, "kEpStatusRemoteReadyTimeout"); - } -} - -void TestDispatchDemoRunsCombine() -{ - const std::string path = "tests/ep/demo/tilexr_ep_dispatch_demo.cpp"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckContains(path, contents, "TileXRMoeEpCombine"); - CheckContains(path, contents, "ValidateCombineOutputs"); - CheckContains(path, contents, "combine validation"); -} - -void TestKernelForwardsActiveMask() -{ - std::string hostLaunch; - if (ReadFile("src/ep/host/ep_kernel_launch.cpp", &hostLaunch)) { - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "xActiveMask"); - } - - std::string kernel; - if (ReadFile("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", &kernel)) { - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "xActiveMaskGM"); - } - - std::string dispatchHelpers; - if (ReadFile("src/ep/kernels/tilexr_ep_dispatch_helpers.h", &dispatchHelpers)) { - CheckContains("src/ep/kernels/tilexr_ep_dispatch_helpers.h", dispatchHelpers, "TileXREpIsTokenActive"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_helpers.h", dispatchHelpers, "xActiveMaskGM == nullptr"); - } -} - -void TestKernelForwardsExpertTokenNumsType() -{ - std::string hostLaunch; - if (ReadFile("src/ep/host/ep_kernel_launch.cpp", &hostLaunch)) { - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "expertTokenNumsType"); - } - - std::string kernel; - if (ReadFile("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", &kernel)) { - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "expertTokenNumsType"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "TileXREpClearExpertTokenNums"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "TileXREpFinalizeExpertTokenNums"); - } - - std::string combineHelpers; - if (ReadFile("src/ep/kernels/tilexr_ep_combine_helpers.h", &combineHelpers)) { - CheckContains("src/ep/kernels/tilexr_ep_combine_helpers.h", combineHelpers, - "TileXREpIncrementExpertTokenNum"); - CheckContains("src/ep/kernels/tilexr_ep_combine_helpers.h", combineHelpers, - "running += expertTokenNumsOut[localExpert]"); - } -} - -void TestKernelForwardsTpRecvCountsOut() -{ - std::string hostLaunch; - if (ReadFile("src/ep/host/ep_kernel_launch.cpp", &hostLaunch)) { - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "tpRecvCountsOut"); - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "tpWorldSize"); - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "tpRankId"); - } - - std::string kernel; - if (ReadFile("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", &kernel)) { - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "tpRecvCountsOutGM"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "TileXREpAppendTpGroupRows"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "TileXREpTpGroupStartRank"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "tpWorldSize"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "tpRankId"); - } - - const std::string demoPath = "tests/ep/demo/tilexr_ep_dispatch_demo.cpp"; - std::string demo; - if (ReadFile(demoPath, &demo)) { - CheckContains(demoPath, demo, "TILEXR_EP_DEMO_TP_RECV_COUNTS"); - CheckContains(demoPath, demo, "TILEXR_EP_DEMO_TP_WORLD_SIZE"); - CheckContains(demoPath, demo, "TILEXR_EP_DEMO_TP_RANK_ID"); - CheckContains(demoPath, demo, "BuildExpectedTpRoutes"); - CheckContains(demoPath, demo, "ValidateTpRecvCounts"); - } -} - -void TestDispatchDemoExercisesV2OptionalInputs() -{ - const std::string path = "tests/ep/demo/tilexr_ep_dispatch_demo.cpp"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckContains(path, contents, "TILEXR_EP_DEMO_ACTIVE_MASK"); - CheckContains(path, contents, "TILEXR_EP_DEMO_EXPERT_TOKEN_NUMS_TYPE"); - CheckContains(path, contents, "xActiveMaskDev"); - CheckContains(path, contents, "expertTokenNumsType"); -} - -void TestKernelForwardsSharedExpertConfig() -{ - std::string hostLaunch; - if (ReadFile("src/ep/host/ep_kernel_launch.cpp", &hostLaunch)) { - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "sharedExpertNum"); - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "sharedExpertRankNum"); - } - - std::string kernel; - if (ReadFile("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", &kernel)) { - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "TileXREpRouteToDstRank"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "sharedExpertRankNum"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "TileXREpLocalExpertCount"); - } - - std::string demo; - if (ReadFile("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", &demo)) { - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "TILEXR_EP_DEMO_SHARED_EXPERT_NUM"); - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "TILEXR_EP_DEMO_SHARED_EXPERT_RANK_NUM"); - } -} - -void TestKernelForwardsStaticQuantConfig() -{ - std::string hostLaunch; - if (ReadFile("src/ep/host/ep_kernel_launch.cpp", &hostLaunch)) { - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "scales"); - CheckContains("src/ep/host/ep_kernel_launch.cpp", hostLaunch, "quantMode"); - } - - std::string kernel; - if (ReadFile("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", &kernel)) { - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "scalesGM"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "quantMode"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "TileXREpCopyStaticQuantRoutePayload"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "TileXREpClampInt8"); - } - - std::string demo; - if (ReadFile("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", &demo)) { - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "TILEXR_EP_DEMO_QUANT_MODE"); - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "TILEXR_EP_DEMO_STATIC_QUANT_SCALE"); - } -} - -void TestKernelForwardsPerTokenDynamicQuantConfig() -{ - std::string kernel; - if (ReadFile("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", &kernel)) { - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "kEpQuantModePerTokenDynamic"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, - "TileXREpCopyPerTokenDynamicQuantRoutePayload"); - CheckContains("src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", kernel, "dynamicScalesOutGM"); - } - - std::string demo; - if (ReadFile("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", &demo)) { - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "usePerTokenDynamicQuant"); - CheckContains("tests/ep/demo/tilexr_ep_dispatch_demo.cpp", demo, "DynamicScaleForXValue"); - } -} - -void TestClearLocalWindowDoesNotPreclearSlotHeaders() -{ - const std::string path = "src/ep/kernels/tilexr_ep_dispatch_kernel.cpp"; - std::string contents; - if (!ReadFile(path, &contents)) { - return; - } - - CheckNotContains(path, contents, - "for (int32_t srcRank = 0; srcRank < rankSize; ++srcRank) {\n" - " StoreSlotHeader(localWindow + SlotOffset(srcRank, slotBytes), 0, srcRank, 0, 0, tBuf);\n" - " }"); -} - -void TestNoForbiddenDependencies() -{ - const std::vector paths = { - "src/ep/kernels/tilexr_ep_dispatch_kernel.cpp", - "src/ep/kernels/tilexr_ep_combine_kernel.cpp", - "src/ep/kernels/tilexr_ep_kernel_common.h", - "src/ep/kernels/tilexr_ep_dispatch_helpers.h", - "src/ep/kernels/tilexr_ep_combine_helpers.h", - "src/ep/host/ep_kernel_launch.cpp", - "src/ep/CMakeLists.txt", - }; - const std::vector forbidden = { - "ops-transformer", - "GetHcclContext", - "TileXRUDMARegister", - "shmem", - }; - - for (std::vector::const_iterator path = paths.begin(); path != paths.end(); ++path) { - std::string contents; - if (!ReadFile(*path, &contents)) { - continue; - } - for (std::vector::const_iterator needle = forbidden.begin(); needle != forbidden.end(); ++needle) { - CheckNotContains(*path, contents, *needle); - } - } -} - -} // namespace - -int main() -{ - TestKernelUsesTileXRPeerMemory(); - TestCrossNodeDispatchUsesUDMARegistry(); - TestCrossNodeDispatchPullsRemoteSlots(); - TestCrossNodeDispatchSeparatesLocalAndRemotePeers(); - TestHostDispatchSplitsCrossNodeKernel(); - TestDispatchHelpersLiveInDispatchHelperFile(); - TestCombineHelpersLiveInCombineHelperFile(); - TestCombineKernelUsesTileXRPeerMemory(); - TestKernelCommonHasCombineHelpers(); - TestDispatchDemoRunsCombine(); - TestKernelForwardsActiveMask(); - TestKernelForwardsExpertTokenNumsType(); - TestKernelForwardsTpRecvCountsOut(); - TestDispatchDemoExercisesV2OptionalInputs(); - TestKernelForwardsSharedExpertConfig(); - TestKernelForwardsStaticQuantConfig(); - TestKernelForwardsPerTokenDynamicQuantConfig(); - TestClearLocalWindowDoesNotPreclearSlotHeaders(); - TestNoForbiddenDependencies(); - if (g_failures != 0) { - std::cerr << g_failures << " TileXR EP kernel source checks failed" << std::endl; - return 1; - } - std::cout << "TileXR EP kernel source checks passed" << std::endl; - return 0; -} diff --git a/tests/memory/CMakeLists.txt b/tests/memory/CMakeLists.txt index bc49e6ad..b636931c 100644 --- a/tests/memory/CMakeLists.txt +++ b/tests/memory/CMakeLists.txt @@ -59,20 +59,7 @@ link_directories( ${ASCEND_HOME_PATH}/${ARCH}-linux/devlib ) -add_executable(test_tilexr_memory_demo_sources - unit/test_tilexr_memory_demo_sources.cpp -) - -target_include_directories(test_tilexr_memory_demo_sources PRIVATE - ${TILEXR_ROOT}/src/include -) -target_compile_definitions(test_tilexr_memory_demo_sources PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - -set(INSTALL_TARGETS - test_tilexr_memory_demo_sources -) +set(INSTALL_TARGETS) if(BUILD_TILEXR_MEMORY_DEMO) find_program(BISHENG_EXECUTABLE bisheng) @@ -191,7 +178,9 @@ if(BUILD_TILEXR_MEMORY_DEMO) endif() endif() -install(TARGETS ${INSTALL_TARGETS} - RUNTIME DESTINATION ${CMAKE_INSTALL_PREFIX}/bin - LIBRARY DESTINATION ${CMAKE_INSTALL_PREFIX}/lib -) +if(INSTALL_TARGETS) + install(TARGETS ${INSTALL_TARGETS} + RUNTIME DESTINATION ${CMAKE_INSTALL_PREFIX}/bin + LIBRARY DESTINATION ${CMAKE_INSTALL_PREFIX}/lib + ) +endif() diff --git a/tests/memory/README.md b/tests/memory/README.md index b9fecff2..231978ee 100644 --- a/tests/memory/README.md +++ b/tests/memory/README.md @@ -23,13 +23,12 @@ cd /path/to/TileXR/tests/memory bash build.sh ``` -If `bisheng` is unavailable, the source-level check still builds while the AICore demo is skipped. +If `bisheng` is unavailable, the AICore demo is skipped. ## Run ```bash cd /path/to/TileXR/tests/memory -./install/bin/test_tilexr_memory_demo_sources bash demo/run_tilexr_memory_demo.sh 2 16 2 0 ``` diff --git a/tests/memory/build.sh b/tests/memory/build.sh index 36bd0007..57388065 100755 --- a/tests/memory/build.sh +++ b/tests/memory/build.sh @@ -34,24 +34,20 @@ fi cmake -DCMAKE_INSTALL_PREFIX="${INSTALL_DIR}" ${DEMO_OPTION} .. make -j$(nproc) -make install +cmake --install . echo "" echo "==========================================" echo " Build Complete" echo "==========================================" -echo "Test binaries installed to: ${INSTALL_DIR}/bin" +echo "Binaries installed to: ${INSTALL_DIR}/bin" echo "" -echo "Available tests:" -echo " - test_tilexr_memory_demo_sources : peer-memory DataCopy demo source checks" if [ -f "${INSTALL_DIR}/bin/tilexr_memory_demo" ]; then echo " - tilexr_memory_demo : TileXR peer-memory DataCopy communication demo" else echo " - tilexr_memory_demo : skipped (requires bisheng/AICore toolchain)" fi echo "" -echo "Run source checks with:" -echo " ./install/bin/test_tilexr_memory_demo_sources" echo "Run demo with:" echo " bash demo/run_tilexr_memory_demo.sh 2 16" echo "==========================================" diff --git a/tests/memory/unit/test_tilexr_memory_demo_sources.cpp b/tests/memory/unit/test_tilexr_memory_demo_sources.cpp deleted file mode 100644 index 52cee587..00000000 --- a/tests/memory/unit/test_tilexr_memory_demo_sources.cpp +++ /dev/null @@ -1,125 +0,0 @@ -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -#define CHECK_TRUE(expr) \ - do { \ - if (!(expr)) { \ - std::cerr << "CHECK_TRUE failed at line " << __LINE__ << ": " #expr << std::endl; \ - ++g_failures; \ - } \ - } while (0) - -std::string RepoPath(const std::string& path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -std::string ReadFile(const std::string& path) -{ - const std::string fullPath = RepoPath(path); - std::ifstream input(fullPath.c_str()); - if (!input.is_open()) { - std::cerr << "failed to open " << fullPath << std::endl; - ++g_failures; - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -void CheckContains(const std::string& path, const std::string& text, const std::string& needle) -{ - if (text.find(needle) == std::string::npos) { - std::cerr << path << " does not contain required text: " << needle << std::endl; - ++g_failures; - } -} - -void CheckNotContains(const std::string& path, const std::string& text, const std::string& needle) -{ - const auto pos = text.find(needle); - if (pos != std::string::npos) { - std::cerr << path << " unexpectedly contains text: " << needle << " at byte " << pos << std::endl; - ++g_failures; - } -} - -void CheckAppearsBefore( - const std::string& path, const std::string& text, const std::string& beforeNeedle, const std::string& afterNeedle) -{ - const auto beforePos = text.find(beforeNeedle); - const auto afterPos = text.find(afterNeedle); - if (beforePos == std::string::npos || afterPos == std::string::npos || beforePos > afterPos) { - std::cerr << path << " expected " << beforeNeedle << " before " << afterNeedle << std::endl; - ++g_failures; - } -} - -void TestMemoryDemoKernelUsesPeerMemorySemantics() -{ - const std::string path = "tests/memory/demo/tilexr_memory_demo_kernel.cpp"; - const std::string text = ReadFile(path); - CheckAppearsBefore(path, text, "#include \"comm_args.h\"", "#include \"kernel_operator.h\""); - CheckContains(path, text, "tilexr_memory_all_gather_kernel"); - CheckContains(path, text, "peerMems"); - CheckContains(path, text, "SyncCollectives"); - CheckContains(path, text, "AscendC::DataCopy"); - CheckContains(path, text, "IPC_DATA_OFFSET"); - CheckNotContains(path, text, "TileXRUDMARegister"); - CheckNotContains(path, text, "UDMAPut"); - CheckNotContains(path, text, "UDMAGet"); -} - -void TestCommArgsOwnsAicoreGmAddrCompatibility() -{ - const std::string path = "src/include/comm_args.h"; - const std::string text = ReadFile(path); - CheckContains(path, text, "defined(__CCE__)"); - CheckContains(path, text, "defined(__CCE_IS_AICORE__)"); - CheckContains(path, text, "TILEXR_ASCENDC_AICORE_COMPILE"); - CheckNotContains(path, text, "__DAV_C220_VEC__"); - CheckNotContains(path, text, "__DAV_C310__"); - CheckAppearsBefore(path, text, "#include \"kernel_operator.h\"", "#ifndef GM_ADDR"); -} - -void TestMemoryDemoHostAndRunnerExist() -{ - const std::string hostPath = "tests/memory/demo/tilexr_memory_demo.cpp"; - const std::string hostText = ReadFile(hostPath); - CheckContains(hostPath, hostText, "TileXRCommInitRankLocal"); - CheckContains(hostPath, hostText, "TileXRGetCommArgsDev"); - CheckContains(hostPath, hostText, "launch_tilexr_memory_all_gather"); - CheckNotContains(hostPath, hostText, "TileXRUDMARegister"); - - const std::string runPath = "tests/memory/demo/run_tilexr_memory_demo.sh"; - const std::string runText = ReadFile(runPath); - CheckContains(runPath, runText, "tilexr_memory_demo"); - CheckContains(runPath, runText, "${TILEXR_ROOT}/install/lib64"); - CheckNotContains(runPath, runText, "/usr/local/lib"); -} - -} // namespace - -int main() -{ - TestMemoryDemoKernelUsesPeerMemorySemantics(); - TestCommArgsOwnsAicoreGmAddrCompatibility(); - TestMemoryDemoHostAndRunnerExist(); - if (g_failures != 0) { - std::cerr << g_failures << " TileXR memory demo source checks failed" << std::endl; - return 1; - } - std::cout << "TileXR memory demo source checks passed" << std::endl; - return 0; -} diff --git a/tests/sdma/CMakeLists.txt b/tests/sdma/CMakeLists.txt index 17c85a30..a56bb6ec 100644 --- a/tests/sdma/CMakeLists.txt +++ b/tests/sdma/CMakeLists.txt @@ -87,20 +87,6 @@ target_include_directories(test_tilexr_sdma_transport_disabled PRIVATE ${ASCEND_DRIVER_PATH}/kernel/inc ) -add_executable(test_tilexr_sdma_comm_wiring - unit/test_tilexr_sdma_comm_wiring.cpp -) -target_compile_definitions(test_tilexr_sdma_comm_wiring PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - -add_executable(test_tilexr_sdma_source_guard - unit/test_tilexr_sdma_source_guard.cpp -) -target_compile_definitions(test_tilexr_sdma_source_guard PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - add_executable(test_tilexr_sdma_header_compile unit/test_tilexr_sdma_header_compile.cpp ) @@ -132,8 +118,6 @@ set(INSTALL_TARGETS test_tilexr_sdma_metadata test_tilexr_sdma_api_invalid test_tilexr_sdma_transport_disabled - test_tilexr_sdma_comm_wiring - test_tilexr_sdma_source_guard test_tilexr_sdma_header_compile test_tilexr_sdma_disabled_comm ) diff --git a/tests/sdma/run_tests.sh b/tests/sdma/run_tests.sh index 793cb9bb..f0035f57 100755 --- a/tests/sdma/run_tests.sh +++ b/tests/sdma/run_tests.sh @@ -63,8 +63,6 @@ fi "${INSTALL_DIR}/bin/test_tilexr_sdma_metadata" "${INSTALL_DIR}/bin/test_tilexr_sdma_transport_disabled" -"${INSTALL_DIR}/bin/test_tilexr_sdma_comm_wiring" -"${INSTALL_DIR}/bin/test_tilexr_sdma_source_guard" "${INSTALL_DIR}/bin/test_tilexr_sdma_header_compile" if [ "${HAL_AVAILABLE}" -eq 1 ]; then diff --git a/tests/sdma/unit/test_tilexr_sdma_comm_wiring.cpp b/tests/sdma/unit/test_tilexr_sdma_comm_wiring.cpp deleted file mode 100644 index cd48b944..00000000 --- a/tests/sdma/unit/test_tilexr_sdma_comm_wiring.cpp +++ /dev/null @@ -1,58 +0,0 @@ -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -std::string RepoPath(const std::string& path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -std::string ReadFile(const std::string& path) -{ - std::ifstream input(RepoPath(path).c_str()); - if (!input.is_open()) { - std::cerr << "failed to open " << RepoPath(path) << std::endl; - ++g_failures; - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -void CheckContains(const std::string& path, const std::string& needle) -{ - const auto text = ReadFile(path); - if (text.find(needle) == std::string::npos) { - std::cerr << "expected text not found in " << path << ": " << needle << std::endl; - ++g_failures; - } -} - -} // namespace - -int main() -{ - CheckContains("src/comm/tilexr_comm.h", "std::unique_ptr sdmaTransport_"); - CheckContains("src/comm/tilexr_comm.cpp", "int TileXRComm::InitSDMA()"); - CheckContains("src/comm/tilexr_comm.cpp", "commArgs_.sdmaWorkspacePtr = sdmaWorkspaceDev_"); - CheckContains("src/comm/tilexr_comm.cpp", "commArgs_.extraFlag |= ExtraFlag::SDMA"); - CheckContains("src/comm/comm_wrap.cpp", "*available = c->IsSDMAAvailable()"); - CheckContains("src/comm/comm_wrap.cpp", "*workspace = c->GetSDMAWorkspacePtr()"); - - if (g_failures != 0) { - std::cerr << g_failures << " SDMA comm wiring checks failed" << std::endl; - return 1; - } - std::cout << "TileXR SDMA comm wiring checks passed" << std::endl; - return 0; -} diff --git a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp deleted file mode 100644 index cfe96e35..00000000 --- a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp +++ /dev/null @@ -1,123 +0,0 @@ -#include -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -std::string RepoPath(const std::string& path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -std::string ReadFile(const std::string& path, bool required = true) -{ - std::ifstream input(RepoPath(path).c_str()); - if (!input.is_open()) { - if (required) { - std::cerr << "failed to open " << RepoPath(path) << std::endl; - ++g_failures; - } - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -void CheckNoNeedle(const std::string& path, const std::string& text, const std::string& needle) -{ - const auto pos = text.find(needle); - if (pos != std::string::npos) { - std::cerr << "unexpected dependency in " << path << ": " << needle - << " at byte " << pos << std::endl; - ++g_failures; - } -} - -void CheckNeedle(const std::string& path, const std::string& text, const std::string& needle) -{ - if (text.find(needle) == std::string::npos) { - std::cerr << "expected text not found in " << path << ": " << needle << std::endl; - ++g_failures; - } -} - -void TestCommSourcesDoNotUseShmem() -{ - const std::vector paths = { - "src/comm/CMakeLists.txt", - "src/comm/tilexr_comm.cpp", - "src/comm/comm_wrap.cpp", - "src/comm/tilexr_comm.h", - "src/comm/sdma/tilexr_sdma_transport.cpp", - "src/comm/sdma/tilexr_sdma_transport.h", - }; - const std::vector forbidden = { - "shmem", - "shmem.h", - "libshmem", - "aclshmem", - "ACLSHMEM", - }; - for (const auto& path : paths) { - const auto text = ReadFile(path); - for (const auto& needle : forbidden) { - CheckNoNeedle(path, text, needle); - } - } -} - -void TestOnlyCompatIncludesSdmaIntrinsics() -{ - const std::vector paths = { - "src/include/tilexr_sdma.h", - "src/include/tilexr_sdma_compat.h", - "src/include/tilexr_sdma_types.h", - "src/include/comm_args.h", - "src/comm/tilexr_comm.cpp", - "src/comm/tilexr_comm.h", - "src/comm/comm_wrap.cpp", - "src/comm/sdma/tilexr_sdma_transport.cpp", - "src/comm/sdma/tilexr_sdma_transport.h", - }; - const std::string ptoPrefix = "pto/npu/comm/async/sdma/"; - const std::string intrinHeader = ptoPrefix + "sdma_async_intrin.hpp"; - const std::string workspaceHeader = ptoPrefix + "sdma_workspace_manager.hpp"; - for (const auto& path : paths) { - const bool required = path != "src/include/tilexr_sdma.h"; - const auto text = ReadFile(path, required); - if (path == "src/include/tilexr_sdma_compat.h") { - CheckNeedle(path, text, intrinHeader); - CheckNoNeedle(path, text, workspaceHeader); - continue; - } - if (path == "src/comm/sdma/tilexr_sdma_transport.cpp") { - CheckNeedle(path, text, workspaceHeader); - CheckNoNeedle(path, text, intrinHeader); - continue; - } - CheckNoNeedle(path, text, ptoPrefix); - } -} - -} // namespace - -int main() -{ - TestCommSourcesDoNotUseShmem(); - TestOnlyCompatIncludesSdmaIntrinsics(); - if (g_failures != 0) { - std::cerr << g_failures << " SDMA source guard checks failed" << std::endl; - return 1; - } - std::cout << "TileXR SDMA source guard checks passed" << std::endl; - return 0; -} diff --git a/tests/udma/CMakeLists.txt b/tests/udma/CMakeLists.txt index 653808c9..6c1e41ab 100644 --- a/tests/udma/CMakeLists.txt +++ b/tests/udma/CMakeLists.txt @@ -83,22 +83,6 @@ target_include_directories(test_tilexr_udma_transport_layout PRIVATE ${TILEXR_ROOT}/src/comm ) -add_executable(test_tilexr_udma_demo_sources - unit/test_tilexr_udma_demo_sources.cpp -) - -target_compile_definitions(test_tilexr_udma_demo_sources PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - -add_executable(test_tilexr_udma_source_guard - unit/test_tilexr_udma_source_guard.cpp -) - -target_compile_definitions(test_tilexr_udma_source_guard PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - # 集成测试：TileXR UDMA add_executable(test_tilexr_udma integration/test_tilexr_udma.cpp @@ -115,8 +99,6 @@ set(INSTALL_TARGETS test_tilexr_udma test_tilexr_udma_registry test_tilexr_udma_transport_layout - test_tilexr_udma_demo_sources - test_tilexr_udma_source_guard ) if(BUILD_TILEXR_UDMA_DEMO) diff --git a/tests/udma/build.sh b/tests/udma/build.sh index 02177a27..af0e90db 100755 --- a/tests/udma/build.sh +++ b/tests/udma/build.sh @@ -62,7 +62,6 @@ echo "" echo "Available tests:" echo " - test_tilexr_udma_transport_layout : UDMA info layout unit tests" echo " - test_tilexr_udma_registry : registered-memory metadata unit tests" -echo " - test_tilexr_udma_source_guard : UDMA ownership/source boundary checks" echo " - test_tilexr_udma : TileXR integration tests" if [ -f "${INSTALL_DIR}/bin/tilexr_udma_demo" ]; then echo " - tilexr_udma_demo : TileXR UDMA communication demo" diff --git a/tests/udma/run_tests.sh b/tests/udma/run_tests.sh index e899d2de..f020b583 100755 --- a/tests/udma/run_tests.sh +++ b/tests/udma/run_tests.sh @@ -46,7 +46,6 @@ fi # 检查测试二进制是否存在 if [ ! -f "${INSTALL_DIR}/bin/test_tilexr_udma_transport_layout" ] || [ ! -f "${INSTALL_DIR}/bin/test_tilexr_udma_registry" ] || - [ ! -f "${INSTALL_DIR}/bin/test_tilexr_udma_source_guard" ] || [ ! -f "${INSTALL_DIR}/bin/test_tilexr_udma" ]; then echo "ERROR: Test binaries not found. Please run build.sh first." exit 1 @@ -68,27 +67,19 @@ echo "==========================================" TEST2_RESULT=$? echo "" -# 测试 3: UDMA source guard（host-only） +# 测试 3: TileXR 集成测试（单进程，单卡） echo "==========================================" -echo "Test 3: TileXR UDMA Source Guard Unit Test" -echo "==========================================" -"${INSTALL_DIR}/bin/test_tilexr_udma_source_guard" -TEST3_RESULT=$? -echo "" - -# 测试 4: TileXR 集成测试（单进程，单卡） -echo "==========================================" -echo "Test 4: TileXR Integration Tests (Single Process)" +echo "Test 3: TileXR Integration Tests (Single Process)" echo "==========================================" export RANK=0 export RANK_SIZE=1 "${INSTALL_DIR}/bin/test_tilexr_udma" -TEST4_RESULT=$? +TEST3_RESULT=$? echo "" -# 测试 5: TileXR 多进程测试（需要 mpirun） +# 测试 4: TileXR 多进程测试（需要 mpirun） echo "==========================================" -echo "Test 5: TileXR Multi-Process Tests (MPI)" +echo "Test 4: TileXR Multi-Process Tests (MPI)" echo "==========================================" # 检查是否有 mpirun @@ -107,14 +98,14 @@ if command -v mpirun &> /dev/null; then unset RANK unset RANK_SIZE mpirun -n 2 "${INSTALL_DIR}/bin/test_tilexr_udma" - TEST5_RESULT=$? + TEST4_RESULT=$? else echo "SKIP: Need at least 2 usable NPUs for multi-rank test" - TEST5_RESULT=0 + TEST4_RESULT=0 fi else echo "SKIP: mpirun not found, skipping multi-process tests" - TEST5_RESULT=0 + TEST4_RESULT=0 fi echo "" @@ -124,14 +115,13 @@ echo " Test Results Summary" echo "==========================================" echo "Test 1 (UDMA Layout): $([ $TEST1_RESULT -eq 0 ] && echo 'PASS' || echo 'FAIL')" echo "Test 2 (UDMA Registry): $([ $TEST2_RESULT -eq 0 ] && echo 'PASS' || echo 'FAIL')" -echo "Test 3 (Source Guard): $([ $TEST3_RESULT -eq 0 ] && echo 'PASS' || echo 'FAIL')" -echo "Test 4 (TileXR Single): $([ $TEST4_RESULT -eq 0 ] && echo 'PASS' || echo 'FAIL')" -echo "Test 5 (TileXR Multi): $([ $TEST5_RESULT -eq 0 ] && echo 'PASS' || echo 'SKIP/FAIL')" +echo "Test 3 (TileXR Single): $([ $TEST3_RESULT -eq 0 ] && echo 'PASS' || echo 'FAIL')" +echo "Test 4 (TileXR Multi): $([ $TEST4_RESULT -eq 0 ] && echo 'PASS' || echo 'SKIP/FAIL')" echo "==========================================" # 返回失败状态 if [ $TEST1_RESULT -ne 0 ] || [ $TEST2_RESULT -ne 0 ] || [ $TEST3_RESULT -ne 0 ] || - [ $TEST4_RESULT -ne 0 ] || [ $TEST5_RESULT -ne 0 ]; then + [ $TEST4_RESULT -ne 0 ]; then exit 1 fi diff --git a/tests/udma/unit/test_tilexr_udma_demo_sources.cpp b/tests/udma/unit/test_tilexr_udma_demo_sources.cpp deleted file mode 100644 index c6ed7f0c..00000000 --- a/tests/udma/unit/test_tilexr_udma_demo_sources.cpp +++ /dev/null @@ -1,90 +0,0 @@ -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -std::string RepoPath(const std::string &path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -std::string ReadFile(const std::string &path) -{ - const std::string fullPath = RepoPath(path); - std::ifstream input(fullPath.c_str()); - if (!input) { - std::cerr << "failed to open " << fullPath << std::endl; - ++g_failures; - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -void CheckContains(const std::string &path, const std::string &text, const std::string &needle) -{ - if (text.find(needle) == std::string::npos) { - std::cerr << path << " missing: " << needle << std::endl; - ++g_failures; - } -} - -void CheckNotContains(const std::string &path, const std::string &text, const std::string &needle) -{ - if (text.find(needle) != std::string::npos) { - std::cerr << path << " should not contain: " << needle << std::endl; - ++g_failures; - } -} - -} // namespace - -int main() -{ - const std::string demoPath = "tests/udma/demo/tilexr_udma_demo.cpp"; - const std::string hccpDefsPath = "src/comm/udma/tilexr_hccp_defs.h"; - const std::string demo = ReadFile(demoPath); - const std::string hccpDefs = ReadFile(hccpDefsPath); - CheckContains(demoPath, demo, "TileXRUDMARegister"); - CheckContains(demoPath, demo, "TileXRUDMAUnregister"); - CheckContains(demoPath, demo, "ExtraFlag::UDMA"); - CheckContains(demoPath, demo, "launch_tilexr_udma_all_gather"); - CheckContains(demoPath, demo, "launch_tilexr_udma_put_signal"); - CheckContains(demoPath, demo, "DemoBarrierAll"); - CheckNotContains(demoPath, demo, "aclshmem"); - CheckNotContains(demoPath, demo, "shmem_"); - - const std::string kernelPath = "tests/udma/demo/tilexr_udma_demo_kernel.cpp"; - const std::string kernel = ReadFile(kernelPath); - CheckContains(kernelPath, kernel, "UDMARegistryEnabled"); - CheckContains(kernelPath, kernel, "UDMAPutNbi"); - CheckContains(kernelPath, kernel, "UDMAPutSignalNbi"); - CheckContains(kernelPath, kernel, "UDMAQuiet"); - CheckNotContains(kernelPath, kernel, "aclshmem"); - CheckNotContains(kernelPath, kernel, "shmem_"); - - CheckContains(hccpDefsPath, hccpDefs, "MEM_SEG_ACCESS_LOCAL_ONLY = 1"); - CheckContains(hccpDefsPath, hccpDefs, "MEM_SEG_ACCESS_READ = (1 << 1)"); - CheckContains(hccpDefsPath, hccpDefs, "MEM_SEG_ACCESS_WRITE = (1 << 2)"); - CheckContains(hccpDefsPath, hccpDefs, "MEM_SEG_ACCESS_ATOMIC = (1 << 3)"); - CheckContains(hccpDefsPath, hccpDefs, - "MEM_SEG_ACCESS_DEFAULT = MEM_SEG_ACCESS_READ | MEM_SEG_ACCESS_WRITE | MEM_SEG_ACCESS_ATOMIC"); - - const std::string transportPath = "src/comm/udma/tilexr_udma_transport.cpp"; - const std::string transport = ReadFile(transportPath); - CheckContains(transportPath, transport, "ResolveLocalEidRoute(rootInfo, topoEdges, localId, allLocalIds[peer], localEid)"); - - if (g_failures == 0) { - std::cout << "TileXR UDMA demo source checks passed" << std::endl; - } - return g_failures == 0 ? 0 : 1; -} diff --git a/tests/udma/unit/test_tilexr_udma_source_guard.cpp b/tests/udma/unit/test_tilexr_udma_source_guard.cpp deleted file mode 100644 index fb9f768b..00000000 --- a/tests/udma/unit/test_tilexr_udma_source_guard.cpp +++ /dev/null @@ -1,198 +0,0 @@ -#include -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -std::string RepoPath(const std::string& path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -std::string ReadFile(const std::string& path) -{ - std::ifstream input(RepoPath(path).c_str()); - if (!input.is_open()) { - std::cerr << "failed to open " << RepoPath(path) << std::endl; - ++g_failures; - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -void CheckContains(const std::string& path, const std::string& text, const std::string& needle) -{ - if (text.find(needle) == std::string::npos) { - std::cerr << "expected text not found in " << path << ": " << needle << std::endl; - ++g_failures; - } -} - -void CheckNotContains(const std::string& path, const std::string& text, const std::string& needle) -{ - const auto pos = text.find(needle); - if (pos != std::string::npos) { - std::cerr << "unexpected text in " << path << ": " << needle - << " at byte " << pos << std::endl; - ++g_failures; - } -} - -void CheckNoNeedles(const std::string& path, const std::vector& needles) -{ - const auto text = ReadFile(path); - for (const auto& needle : needles) { - CheckNotContains(path, text, needle); - } -} - -void TestTileXRCommUsesUDMAContextBoundary() -{ - const std::string headerPath = "src/comm/tilexr_comm.h"; - const auto header = ReadFile(headerPath); - - CheckContains(headerPath, header, "class TileXRUDMAContext;"); - CheckContains(headerPath, header, "std::unique_ptr udmaContext_"); - - const std::vector forbiddenHeaderText = { - "class TileXRUDMATransport;", - "GM_ADDR udmaInfoDev_", - "GM_ADDR udmaRegistryDev_", - "GM_ADDR udmaRegisteredPtr_", - "TileXRUDMARegistry udmaRegistry_", - "std::unique_ptr", - }; - for (const auto& needle : forbiddenHeaderText) { - CheckNotContains(headerPath, header, needle); - } -} - -void TestUDMATransportStaysBehindContext() -{ - CheckNoNeedles("src/comm/tilexr_comm.cpp", { - "#include \"udma/tilexr_udma_transport.h\"", - "new (nothrow) TileXRUDMATransport", - "udmaTransport_->", - "udmaTransport_.", - }); - - CheckContains("src/comm/udma/tilexr_udma_context.h", - ReadFile("src/comm/udma/tilexr_udma_context.h"), - "class TileXRUDMAContext"); - CheckContains("src/comm/udma/tilexr_udma_context.cpp", - ReadFile("src/comm/udma/tilexr_udma_context.cpp"), - "#include \"udma/tilexr_udma_transport.h\""); -} - -void TestUDMAContextShutdownIsLocalOnly() -{ - const std::string path = "src/comm/udma/tilexr_udma_context.cpp"; - const auto text = ReadFile(path); - const auto shutdownPos = text.find("void TileXRUDMAContext::Shutdown()"); - const auto registerPos = text.find("int TileXRUDMAContext::RegisterMemory", shutdownPos); - if (shutdownPos == std::string::npos || registerPos == std::string::npos) { - std::cerr << "failed to locate TileXRUDMAContext::Shutdown body" << std::endl; - ++g_failures; - return; - } - - const auto shutdownBody = text.substr(shutdownPos, registerPos - shutdownPos); - CheckNotContains(path, shutdownBody, "UnregisterMemory("); - CheckContains(path, shutdownBody, "transport_->Shutdown();"); -} - -void TestUDMAReviewFeedbackGuards() -{ - const std::string contextPath = "src/comm/udma/tilexr_udma_context.cpp"; - const auto context = ReadFile(contextPath); - CheckContains(contextPath, context, - "TILEXR_LOG(ERROR) << \"TileXRUDMARegister called while UDMA is unavailable\""); - CheckContains(contextPath, context, - "TILEXR_LOG(ERROR) << \"TileXRUDMARegister is not supported in InitThread mode\""); - CheckContains(contextPath, context, "RollbackTransportRegistration("); - CheckContains(contextPath, context, - "TILEXR_LOG(ERROR) << \"TileXRUDMAUnregister failed to clear comm args: \""); - CheckContains(contextPath, context, - "TILEXR_LOG(ERROR) << \"TileXR UDMA memory unregistration failed: \""); - CheckContains(contextPath, context, - "TILEXR_LOG(ERROR) << \"TileXR UDMA restore registration called with invalid state\""); - CheckContains(contextPath, context, - "TILEXR_LOG(ERROR) << \"TileXR UDMA failed to restore previous registration: \""); - CheckContains(contextPath, context, "TILEXR_LOG(ERROR) << \"Free UDMA registry failed: \""); - - const std::string commPath = "src/comm/tilexr_comm.cpp"; - const auto comm = ReadFile(commPath); - CheckContains(commPath, comm, "TILEXR_LOG(ERROR) << \"TileXR UDMA update comm args failed: \""); - CheckContains(commPath, comm, - "TILEXR_LOG(ERROR) << \"ApplyUDMACommArgsStateCallback missing user data\""); - CheckContains(commPath, comm, - "TILEXR_LOG(ERROR) << \"TileXRUDMARegister called while UDMA is unavailable\""); -} - -void TestPublicHeadersDoNotExposeUDMAContext() -{ - const std::vector publicHeaders = { - "src/include/tilexr_api.h", - "src/include/tilexr_udma.h", - "src/include/tilexr_udma_reg.h", - "src/include/tilexr_udma_types.h", - "src/include/comm_args.h", - }; - for (const auto& path : publicHeaders) { - const auto text = ReadFile(path); - CheckNotContains(path, text, "tilexr_udma_context.h"); - CheckNotContains(path, text, "TileXRUDMAContext"); - } -} - -void TestCommSourcesDoNotUseShmem() -{ - const std::vector paths = { - "src/comm/CMakeLists.txt", - "src/comm/tilexr_comm.cpp", - "src/comm/comm_wrap.cpp", - "src/comm/tilexr_comm.h", - "src/comm/udma/tilexr_udma_context.cpp", - "src/comm/udma/tilexr_udma_context.h", - "src/comm/udma/tilexr_udma_transport.cpp", - "src/comm/udma/tilexr_udma_transport.h", - }; - const std::vector forbidden = { - "shmem", - "shmem.h", - "libshmem", - "aclshmem", - "ACLSHMEM", - }; - for (const auto& path : paths) { - CheckNoNeedles(path, forbidden); - } -} - -} // namespace - -int main() -{ - TestTileXRCommUsesUDMAContextBoundary(); - TestUDMATransportStaysBehindContext(); - TestUDMAContextShutdownIsLocalOnly(); - TestUDMAReviewFeedbackGuards(); - TestPublicHeadersDoNotExposeUDMAContext(); - TestCommSourcesDoNotUseShmem(); - if (g_failures != 0) { - std::cerr << g_failures << " UDMA source guard checks failed" << std::endl; - return 1; - } - std::cout << "TileXR UDMA source guard checks passed" << std::endl; - return 0; -} diff --git a/tools/collectives/run_collective_perf_multihost_suite.sh b/tools/collectives/run_collective_perf_multihost_suite.sh index 007c9ac3..b41760d2 100755 --- a/tools/collectives/run_collective_perf_multihost_suite.sh +++ b/tools/collectives/run_collective_perf_multihost_suite.sh @@ -6,7 +6,7 @@ usage() { Usage: run_collective_perf_multihost_suite.sh prepare run_collective_perf_multihost_suite.sh build - run_collective_perf_multihost_suite.sh guards + run_collective_perf_multihost_suite.sh preflight run_collective_perf_multihost_suite.sh case run_collective_perf_multihost_suite.sh suite run_collective_perf_multihost_suite.sh profile-probe @@ -17,7 +17,7 @@ Environment: TILEXR_PROFILE_BUILD_DIR Profile build dir. Defaults to /build-profile-950. TILEXR_PROFILE_DIR Multi-host profile output root. Defaults to /run/prof/collectives-2host. TILEXR_PROFILE_BUILD_JOBS Build parallelism. Defaults to nproc or 8. - TILEXR_PROFILE_ENV_SCRIPT Optional environment script sourced before build/guards. + TILEXR_PROFILE_ENV_SCRIPT Optional environment script sourced before build/preflight. CMAKE Optional cmake executable name. Defaults to cmake from PATH. Multi-host runs also require TILEXR_MULTIHOST_PEERS and usually TILEXR_COMM_ID. @@ -60,16 +60,12 @@ build_profile() { -DTILEXR_COLLECTIVES_ENABLE_PROFILING=ON \ -DBUILD_TESTING=OFF "${cmake_bin}" --build "${build_dir}" --target \ - test_tilexr_collectives_kernel_ownership \ - test_tilexr_collectives_tools_sources \ tilexr_collective_perf -j"${build_jobs}" } -run_guards() { +run_preflight() { cd "${repo_dir}" source_profile_env - "${build_dir}/tests/collectives/test_tilexr_collectives_kernel_ownership" - "${build_dir}/tests/collectives/test_tilexr_collectives_tools_sources" python3 "${repo_dir}/tests/collectives/unit/test_collective_profile_report.py" } @@ -152,12 +148,12 @@ case "${command}" in build) build_profile ;; - guards) - run_guards + preflight) + run_preflight ;; prepare) build_profile - run_guards + run_preflight ;; case) if [[ $# -ne 4 ]]; then