From ef460bf3c08faabde7657d2323f2e8a3c50a4694 Mon Sep 17 00:00:00 2001 From: zjgarvey Date: Tue, 9 Jun 2026 12:24:02 -0700 Subject: [PATCH] test: sustained H2D/D2H copy integrity 32 iterations of H2D-then-D2H copies of a 4 MiB buffer, each verified with a full element-wise comparison. Exercises sustained host<->device transfer and the stream queue-submission path; any dropped or reordered byte fails. GPU required. Verified on MI300X (98 assertions). Co-Authored-By: Claude Opus 4.8 (1M context) Signed-off-by: zjgarvey --- tests/CMakeLists.txt | 5 ++- tests/aql_doorbell/CMakeLists.txt | 9 +++++ tests/aql_doorbell/aql_doorbell_test.cpp | 47 ++++++++++++++++++++++++ 3 files changed, 60 insertions(+), 1 deletion(-) create mode 100644 tests/aql_doorbell/CMakeLists.txt create mode 100644 tests/aql_doorbell/aql_doorbell_test.cpp diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index f632a71..c28792a 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -24,4 +24,7 @@ add_subdirectory(kernel_smoke) add_subdirectory(multithread) # Graph API: -add_subdirectory(graph) \ No newline at end of file +add_subdirectory(graph) + +# Nightly ROCm bring-up regression tests: +add_subdirectory(aql_doorbell) \ No newline at end of file diff --git a/tests/aql_doorbell/CMakeLists.txt b/tests/aql_doorbell/CMakeLists.txt new file mode 100644 index 0000000..2746406 --- /dev/null +++ b/tests/aql_doorbell/CMakeLists.txt @@ -0,0 +1,9 @@ +# Copyright 2026 The IREE Authors +# +# Licensed under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +add_hip_cts_test(aql_doorbell_test + SOURCES aql_doorbell_test.cpp +) diff --git a/tests/aql_doorbell/aql_doorbell_test.cpp b/tests/aql_doorbell/aql_doorbell_test.cpp new file mode 100644 index 0000000..cd1d3e3 --- /dev/null +++ b/tests/aql_doorbell/aql_doorbell_test.cpp @@ -0,0 +1,47 @@ +// Copyright 2026 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +// Sustained host<->device copy integrity. Many back-to-back H2D/D2H roundtrips +// of a multi-MiB buffer, each followed by a full element-wise comparison, so a +// byte dropped or reordered on any transfer fails the test. Each hipMemcpy also +// submits to the stream's queue, so the loop doubles as a stress test of the +// queue-submission path under sustained use. GPU required. + +#include +#include +#include + +#include +#include "hip_loader.hpp" +#include "hip_test_fixture.hpp" + +TEST_CASE_METHOD(HipTestFixture, + "Repeated H2D/D2H roundtrips preserve data integrity", + "[memcpy][doorbell][nightly]") { + const size_t kCount = 1u << 20; // 1M uint32 = 4 MiB + const size_t kBytes = kCount * sizeof(uint32_t); + std::vector host_in(kCount); + std::vector host_out(kCount, 0u); + for (size_t i = 0; i < kCount; ++i) { + host_in[i] = static_cast(i * 2654435761u); + } + + void* dev = nullptr; + REQUIRE(hip().hipMalloc(&dev, kBytes) == hipSuccess); + + // 32 iterations, each an H2D copy followed by a D2H copy of the same buffer. + // A conformant runtime returns the bytes intact on every round trip. + for (int iter = 0; iter < 32; ++iter) { + REQUIRE(hip().hipMemcpy(dev, host_in.data(), kBytes, + hipMemcpyHostToDevice) == hipSuccess); + std::fill(host_out.begin(), host_out.end(), 0u); + REQUIRE(hip().hipMemcpy(host_out.data(), dev, kBytes, + hipMemcpyDeviceToHost) == hipSuccess); + REQUIRE(host_out == host_in); + } + + REQUIRE(hip().hipFree(dev) == hipSuccess); +}