fix(hrr): replay pitched host copies over exactly the copied rect

projects/clr/hipamd/src/hrr/hip_capture.cpp

@@ -1554,6 +1554,12 @@ static size_t memcpy3d_byte_count(const struct hipMemcpy3DParms* p) {
   return p->extent.width * p->extent.height * p->extent.depth;
 }
 
+// Defined with the driver-copy helpers below. The runtime widens a hipMemcpy3D
+// to the same HIP_MEMCPY3D, so its host side has the same footprint.
+static size_t drvmemcpy_host_byte_count(size_t pitch, size_t pitch_height,
+                                        size_t x, size_t y, size_t z,
+                                        size_t width, size_t height, size_t depth);
+
 // Helper shared by all four 3D variants.
 // Writes H2D blob (src host data) and D2H expected blob (dst host data after copy),
 // then emits the event record.
@@ -1568,14 +1574,26 @@ static void capture_memcpy3d_impl(
   std::memcpy(a.parms_bytes, p, sizeof(a.parms_bytes));
   size_t byte_count = memcpy3d_byte_count(p);
 
-  if (p->kind == hipMemcpyHostToDevice && p->srcPtr.ptr && byte_count > 0) {
-    // H2D: host source is valid at call time — no stream sync needed.
-    auto h = hrr_cap::writer::write_blob(p->srcPtr.ptr, byte_count);
+  if (p->kind == hipMemcpyHostToDevice && p->srcPtr.ptr && byte_count > 0 &&
+      a.ret == hipSuccess) {
+    // H2D: host source is valid at call time, so no stream sync is needed.
+    // Replay reads the blob with the recorded pitch and position, so it spans
+    // srcPtr.ptr through the last copied byte, as the driver copies do.
+    auto h = hrr_cap::writer::write_blob(
+        p->srcPtr.ptr,
+        drvmemcpy_host_byte_count(p->srcPtr.pitch, p->srcPtr.ysize, p->srcPos.x, p->srcPos.y,
+                                  p->srcPos.z, p->extent.width, p->extent.height,
+                                  p->extent.depth));
     a.blob_hash_lo = h.lo;
     a.blob_hash_hi = h.hi;
-  } else if (p->kind == hipMemcpyDeviceToHost && p->dstPtr.ptr && byte_count > 0) {
+  } else if (p->kind == hipMemcpyDeviceToHost && p->dstPtr.ptr && byte_count > 0 &&
+             a.ret == hipSuccess) {
     // D2H: real call already completed (sync API) or stream sync done below;
     // host buffer now holds GPU result — capture it as the expected output.
+    // The blob spans dstPtr.ptr through the last copied byte, laid out with the
+    // recorded pitch and position, so replay can compare exactly the copied
+    // rows. A rejected copy is skipped: nothing validated its rect, and that
+    // span can reach past the caller's buffer.
     if (is_async && stream) {
       hipError_t sync_r = g_real_table.hipStreamSynchronize_fn(stream);
       if (sync_r != hipSuccess) {
@@ -1585,7 +1603,11 @@ static void capture_memcpy3d_impl(
         return;
       }
     }
-    auto h = hrr_cap::writer::write_blob(p->dstPtr.ptr, byte_count);
+    auto h = hrr_cap::writer::write_blob(
+        p->dstPtr.ptr,
+        drvmemcpy_host_byte_count(p->dstPtr.pitch, p->dstPtr.ysize, p->dstPos.x, p->dstPos.y,
+                                  p->dstPos.z, p->extent.width, p->extent.height,
+                                  p->extent.depth));
     a.d2h_hash_lo = h.lo;
     a.d2h_hash_hi = h.hi;
   }
@@ -1650,6 +1672,11 @@ static size_t memcpy2d_host_byte_count(size_t pitch, size_t width, size_t height
 // end whenever pitch > width, height > 1 or an offset is non-zero. The volume is
 // only correct for a fully dense rect, which is also why this never shrinks a
 // blob: for a copy the runtime accepted, footprint >= width*height*depth.
+//
+// D2H expected outputs use the same footprint of dstHost, so replay can re-run
+// the copy into a buffer laid out the same way and compare exactly the copied
+// rows. An older archive holds the flat volume there, which replay tells apart
+// by its size.
 static size_t drvmemcpy_host_byte_count(size_t pitch, size_t pitch_height,
                                         size_t x, size_t y, size_t z,
                                         size_t width, size_t height, size_t depth) {
@@ -1677,13 +1704,12 @@ static void capture_drvmemcpy3d_impl(T& a, hrr_api_id_t api_id,
     }
   } else if (p->dstMemoryType == hipMemoryTypeHost && p->dstHost &&
              p->srcMemoryType != hipMemoryTypeArray) {
-    // D2H expected output. Replay does a flat readback of the copied volume
-    // (it never substitutes dstHost), so the blob stays the flat volume to keep
-    // the two sides the same shape. A pitched destination rect is therefore not
-    // validated faithfully; that is a fidelity gap, not a replay over-read.
+    // D2H expected output, the host footprint of the destination rect.
     // An array source is skipped because playback declines array-typed rects,
     // so the blob would be an expected output nothing ever validates.
-    size_t n = p->WidthInBytes * p->Height * p->Depth;
+    size_t n = drvmemcpy_host_byte_count(p->dstPitch, p->dstHeight, p->dstXInBytes,
+                                         p->dstY, p->dstZ, p->WidthInBytes,
+                                         p->Height, p->Depth);
     if (n > 0) {
       // The null stream needs the sync too: hipDrvMemcpy3DAsync(p, nullptr) is
       // still asynchronous, so skipping it can snapshot dstHost before the copy
@@ -1724,7 +1750,10 @@ static void capture_drvmemcpy2d_impl(T& a, hrr_api_id_t api_id, const hip_Memcpy
     }
   } else if (p->dstMemoryType == hipMemoryTypeHost && p->dstHost &&
              p->srcMemoryType != hipMemoryTypeArray) {
-    size_t n = p->WidthInBytes * p->Height;  // flat volume; see the 3D note above
+    size_t pitch = p->dstPitch ? p->dstPitch : p->dstXInBytes + p->WidthInBytes;
+    size_t n = drvmemcpy_host_byte_count(pitch, /*pitch_height=*/0, p->dstXInBytes,
+                                         p->dstY, /*z=*/0, p->WidthInBytes,
+                                         p->Height, /*depth=*/1);
     if (n > 0) {
       auto h = hrr_cap::writer::write_blob(p->dstHost, n);
       a.d2h_hash_lo = h.lo; a.d2h_hash_hi = h.hi;

projects/hrr/DESIGN.md

@@ -1353,6 +1353,8 @@ The event wire format (finding H5):
   and copied into the translated device `dst` with the recorded `dpitch`. D2H snapshots
   the host `dst` after the copy and validates the device result against it at replay.
   Row-padded image/tensor buffers are now handled.
+- The D2H expected output of `hipMemcpy2D`, the four `hipMemcpy3D` spellings, `hipDrvMemcpy3D` / `hipDrvMemcpy3DAsync` and the three driver 2D spellings is the host destination from its base pointer through the last copied byte, laid out with the recorded pitch and offsets: the same footprint as the driver copies' H2D source. A footprint blob therefore also holds the host bytes between rows and before the first copied row, which the copy never touched. Replay re-runs the recorded copy into a scratch buffer of that size and compares only the copied rows, so bytes between rows or before the first one are never judged. Archives captured earlier hold the flat `width*height*depth` bytes for the 3D and driver copies; replay counts such a check as skipped unless the copy is dense from the base, in which case the two layouts coincide. A recorded rect whose footprint overflows `size_t` is counted as skipped too, so an archive whose every D2H check is skipped fails the replay rather than passing as one with no validation blobs.
+- The H2D source of the four `hipMemcpy3D` spellings is recorded with the same footprint, from `srcPtr.ptr` through the last copied byte, and only for a copy the runtime accepted. Replay substitutes it only when the blob spans the recorded source rect. Otherwise (no blob, or the flat `width*height*depth` blob of an earlier archive) the copy is skipped, rather than read past the blob or issued from the capture-time host address.
 - `hipMemset3D` / `hipMemset3DAsync` drop the destination pitched pointer/extent at
   capture (`pitchedDevPtr = 0`) and no-op at replay, so 3D-memset-initialized regions
   are invisible to replay.
@@ -1372,9 +1374,7 @@ dispatch before the create populates the translation map and silently
 
 D2H validation can pass when replay actually diverged:
 
-- **Length clamp.** Comparison uses `min(copy_size, blob_size)`; a truncated or
-  crash-recovered blob validates only a prefix (the corrupted tail is unchecked) and
-  still counts as PASS. A zero-length compare counts as pass.
+- **Length clamp.** Linear copies compare `min(copy_size, blob_size)`; a truncated or crash-recovered blob validates only a prefix (the corrupted tail is unchecked) and still counts as PASS. A zero-length compare counts as pass. The 2D, 3D and driver copies do not clamp: a blob shorter than the rect's footprint is not validated.
 - **Float-dtype guessing.** Blobs carry no dtype. On a byte mismatch the validator
   tries `{fp32, bf16, fp16, fp64}` and passes on the first encoding within tolerance,
   so integer/index/pointer output buffers can silently false-pass; both-NaN counts as