test: kind-aware AQL doorbell (H2D/D2H roundtrip)

32 H2D/D2H roundtrips with data-integrity checks; each rings the AQL queue
doorbell, so on a USER-kind-doorbell runtime (nightly MI300X) this repeatedly
exercises the hsa_signal_store_screlease doorbell path. Verified on MI300X
(98 assertions).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: zjgarvey <zjgarvey@gmail.com>

Author: zjgarvey

tests/CMakeLists.txt

@@ -24,4 +24,7 @@ add_subdirectory(kernel_smoke)
 add_subdirectory(multithread)
 
 # Graph API:
-add_subdirectory(graph)
+add_subdirectory(graph)
+
+# Nightly ROCm bring-up regression tests:
+add_subdirectory(aql_doorbell)

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
+)

tests/aql_doorbell/aql_doorbell_test.cpp

@@ -0,0 +1,52 @@
+// 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
+
+// Nightly ROCm on MI300X returns a USER-kind queue doorbell signal (not
+// DOORBELL-kind), so the cached doorbell union slot is a signal value, not a
+// writable MMIO address. A raw atomic store to that slot faults — every H2D
+// copy / kernel dispatch would SEGV; the doorbell must be rung via
+// hsa_signal_store_screlease for USER/interrupt-kind doorbells.
+//
+// Each AQL submission rings the doorbell, so a stream of copies with integrity
+// checks exercises that path repeatedly on the affected hardware. GPU required;
+// the SEGV only reproduces on a USER-kind-doorbell runtime.
+
+#include <algorithm>
+#include <cstdint>
+#include <vector>
+
+#include <catch2/catch_test_macros.hpp>
+#include "hip_loader.hpp"
+#include "hip_test_fixture.hpp"
+
+TEST_CASE_METHOD(HipTestFixture,
+                 "Repeated H2D/D2H roundtrips ring the doorbell correctly",
+                 "[memcpy][doorbell][nightly]") {
+  const size_t kCount = 1u << 20;  // 1M uint32 = 4 MiB
+  const size_t kBytes = kCount * sizeof(uint32_t);
+  std::vector<uint32_t> host_in(kCount);
+  std::vector<uint32_t> host_out(kCount, 0u);
+  for (size_t i = 0; i < kCount; ++i) {
+    host_in[i] = static_cast<uint32_t>(i * 2654435761u);
+  }
+
+  void* dev = nullptr;
+  REQUIRE(hip().hipMalloc(&dev, kBytes) == hipSuccess);
+
+  // 32 doorbell rings (H2D + D2H each). On USER-kind-doorbell hardware a raw
+  // doorbell store faults on the first ring; a correct runtime copies and the
+  // data round-trips intact.
+  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);
+}