danielbolsson · Sep 9, 2026 · Aug 20, 2026 · Aug 21, 2026 · Aug 21, 2026 · Aug 24, 2026
@@ -15,7 +15,8 @@
  *   - embedded device pointers inside a by-value struct kernel argument,
  *   - an uncaptured host write that deterministically diverges replay,
  *   - the "null optional output pointer + 0x20000" GPU-fault class, and
- *   - replay zero-init of an otherwise-uninitialised device allocation.
+ *   - replay zero-init of an allocation whose capture-time contents come from
+ *     an unrecorded write (host or HSA runtime).
  *
  * Like hrr_workload.cc, every TEST_CASE here is hidden with the Catch2 [.] tag
  * so it is NOT auto-discovered by CTest.  Each is driven from hrr_roundtrip.cc:
@@ -28,6 +29,10 @@
 
 #include "hrr_test_common.hh"
 
+#ifndef _WIN32
+#include <hsa/hsa_ext_amd.h>
+#endif
+
 #include <cstdint>
 
 // ---------------------------------------------------------------------------
@@ -274,19 +279,30 @@ TEST_CASE("Unit_HRR_NullOptionalPtr_Direct", "[.][hrr-direct]") {
 }
 
 // ===========================================================================
-// A5. Zero-init read of an uninitialised device allocation
+// A5. Zero-init read of an allocation initialized by an unrecorded write
 //
-// hipMalloc a buffer and do NOT initialise it; a kernel copies it to `out`,
-// then D2H.  Fresh ROCm allocations are zeroed, so at capture out == 0 and the
-// recorded blob is all-zero.  Replay with HIP_HRR_REPLAY_ZERO_INIT=1 reproduces
-// the zeroed source deterministically, so the replayed out matches; with the
-// knob off, replay may reuse stale bytes and diverge.
+// hipMalloc a buffer and zero it through the lower-level HSA runtime. HRR
+// records HIP calls only, so replay sees the allocation but not the HSA fill and
+// must reproduce the zeros through HIP_HRR_REPLAY_ZERO_INIT. This makes the
+// capture oracle deterministic without recording a hipMemset that playback
+// would execute independently of the zero-init knob.
 // ===========================================================================
 TEST_CASE("Unit_HRR_ZeroInitRead_Direct", "[.][hrr-direct]") {
   HRR_HIP_CHECK(hipSetDevice(0));
 
-  float* dsrc = nullptr;  // intentionally never written
+  float* dsrc = nullptr;
+#ifndef _WIN32
   HRR_HIP_CHECK(hipMalloc(&dsrc, kSZ));
+  // hsa_amd_memory_fill() counts in 32-bit words, not elements.  kN happens to
+  // be the same number here only because kSZ is float[kN]; spell out the unit
+  // so a change of element type cannot silently under- or over-fill.
+  REQUIRE(hsa_amd_memory_fill(dsrc, 0, kSZ / sizeof(uint32_t)) == HSA_STATUS_SUCCESS);
+#else
+  // Native Windows has no ROCr/HSA runtime. Keep its deterministic oracle via
+  // host stores; Linux retains coverage of the ordinary hipMalloc replay path.
+  HRR_HIP_CHECK(hipMallocManaged(&dsrc, kSZ));
+  for (int i = 0; i < kN; ++i) dsrc[i] = 0.0f;
+#endif
   float* dout = nullptr;
   HRR_HIP_CHECK(hipMalloc(&dout, kSZ));
 

@@ -722,8 +722,9 @@ HRR_TEST_CASE(Unit_HRR_EmbeddedPtrRoundtrip) {
 /**
  * Test Description
  * ----------------
- *   - Capture Unit_HRR_ZeroInitRead_Direct (uninitialised device buffer copied
- *     to out, then D2H; captured blob is all-zero on a zeroing driver).
+ *   - Capture Unit_HRR_ZeroInitRead_Direct (hipMalloc allocation zeroed through
+ *     an unrecorded HSA fill on ROCr, copied to out, then D2H; native Windows
+ *     uses managed memory and unrecorded CPU stores because HSA is unavailable).
  *   - Replay with HIP_HRR_REPLAY_ZERO_INIT=1: the replayed source is zeroed
  *     deterministically, so the D2H validates (exit 0, pass >= 1).  Verifies the
  *     zero-init replay knob.  Note: with the knob off the replay may reuse stale
@@ -741,12 +742,8 @@ HRR_TEST_CASE(Unit_HRR_ZeroInitRoundtrip) {
   int d2h_pass = 0, d2h_fail = 0;
   REQUIRE(hrr_parse_d2h_summary(out, d2h_pass, d2h_fail));
   INFO("D2H pass=" << d2h_pass << " fail=" << d2h_fail);
-// With ASAN enabled this won't be true, because inside Unit_HRR_ZeroInitRead_Direct
-// fresh device allocations are not zeroed with ASAN enabled
-#if !defined(ENABLE_ADDRESS_SANITIZER)
   CHECK(d2h_pass >= 1);
   CHECK(d2h_fail == 0);
-#endif
 }
 
 /**

@@ -2784,6 +2784,12 @@ TEST_CASE("Unit_HRR_DrvMemcpy3D_Direct", "[.][hrr][direct]") {
     p.srcPitch = SZ; p.srcHeight = 1;
     p.dstPitch = SZ; p.dstHeight = 1;
     HRR_HIP_CHECK(hipDrvMemcpy3D(&p)); }
+  // hipDrvMemcpy3D is host-synchronous only when the copy touches host memory;
+  // ihipMemcpyParam3D() forces the D2D case asynchronous, so (3) is merely
+  // enqueued on the null stream.  s was created with hipStreamNonBlocking and so
+  // never joins the null stream, leaving the readback below unordered against
+  // the copy unless it is drained here.
+  HRR_HIP_CHECK(hipDeviceSynchronize());
 
   int* h = new int[N]();
   HRR_HIP_CHECK(hipMemcpyAsync(h, C, SZ, hipMemcpyDeviceToHost, s));
@@ -2894,6 +2900,10 @@ TEST_CASE("Unit_HRR_DrvMemcpy2DUnaligned_Direct", "[.][hrr][direct]") {
     p.dstMemoryType = hipMemoryTypeDevice; p.dstDevice = reinterpret_cast<hipDeviceptr_t>(B); p.dstPitch = SZ;
     p.WidthInBytes  = SZ; p.Height = 1;
     HRR_HIP_CHECK(hipDrvMemcpy2DUnaligned(&p)); }
+  // Same ordering hazard as the 3D workload above: hipDrvMemcpy2DUnaligned goes
+  // through ihipMemcpyParam3D(), which forces the D2D case asynchronous on the
+  // null stream, and s is non-blocking.
+  HRR_HIP_CHECK(hipDeviceSynchronize());
 
   int* h = new int[N]();
   HRR_HIP_CHECK(hipMemcpyAsync(h, B, SZ, hipMemcpyDeviceToHost, s));