feat(hrr): refuse the calls replay cannot reproduce, by name

projects/clr/hipamd/src/hip_code_object.cpp

@@ -708,4 +708,31 @@ void StatCO::ForEachFatBinaryBlob(void (*cb)(const void*)) const {
     cb(data);
   }
 }
+
+// ================================================================================================
+void StatCO::ForEachGlobalVar(void (*cb)(const void*, const char*, size_t,
+                                         const void*)) {
+  std::scoped_lock lock(sclock_);
+  // Resolved here rather than by the caller: this sweep runs from HRR's
+  // capture init, which is inside hip::init(), and re-entering the runtime
+  // through hipGetSymbolAddress from there deadlocks on the init once-flag.
+  // No device is current that early, and GetGlobalVar asserts rather than
+  // erroring on a negative device id, so the address is reported as
+  // unresolved instead.
+  // No device is current this early in init, and GetGlobalVar asserts rather
+  // than erroring on a negative id, so fall back to device 0 — which is the
+  // device a recording resolves its globals on anyway.
+  int device_id = ihipGetDevice();
+  if (device_id < 0) device_id = 0;
+  const bool resolvable = static_cast<size_t>(device_id) < g_devices.size();
+  for (const auto& [host_var, var] : vars_) {
+    if (var == nullptr) continue;
+    hipDeviceptr_t dev_ptr = nullptr;
+    size_t sym_size = 0;
+    if (resolvable &&
+        GetGlobalVar(host_var, device_id, &dev_ptr, &sym_size) != hipSuccess)
+      dev_ptr = nullptr;
+    cb(host_var, var->GetName().c_str(), var->GetSize(), dev_ptr);
+  }
+}
 }  // namespace hip

projects/clr/hipamd/src/hip_code_object.hpp

@@ -193,6 +193,12 @@ class StatCO : public CodeObject {
   // Iterate all registered fat binary data pointers — for HRR capture post-registration sweep.
   void ForEachFatBinaryBlob(void (*cb)(const void*)) const;
 
+  // Iterate all registered __device__ globals as (host shadow address, symbol
+  // name, size, device address) — the same post-registration sweep for HRR,
+  // which otherwise never sees __hipRegisterVar because it fires at
+  // static-init time. The device address is null if it cannot be resolved yet.
+  void ForEachGlobalVar(void (*cb)(const void*, const char*, size_t, const void*));
+
  private:
   mutable std::recursive_mutex sclock_;    //!< Guards Static Code object
   const PlatformState& owner_;             //!< Reference to owning PlatformState

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

@@ -36,6 +36,7 @@
 #include <cerrno>
 #include <filesystem>
 #include <mutex>
+#include <set>
 #include <string>
 #include <unordered_set>
 #include <algorithm>
@@ -216,6 +217,12 @@ static std::atomic<uint64_t> g_blob_count{0};
 static std::mutex                      g_blob_mu;
 static std::unordered_set<std::string> g_written_blobs;
 
+// APIs recorded in this archive that replay cannot reproduce (note_unreplayable).
+// Listed in manifest.json so the gap is a property of the archive rather than
+// something only visible in a replay log.
+static std::mutex                      g_unreplayable_mu;
+static std::set<std::string>           g_unreplayable_apis;
+
 // ---------------------------------------------------------------------------
 // Low-level fd helpers
 // ---------------------------------------------------------------------------
@@ -499,6 +506,18 @@ static void write_manifest_stdio(const char* output_dir, bool complete) {
           complete ? "true" : "false",
           static_cast<unsigned long long>(g_event_count.load()),
           static_cast<unsigned long long>(g_blob_count.load()));
+  {
+    std::lock_guard<std::mutex> lk(g_unreplayable_mu);
+    if (!g_unreplayable_apis.empty()) {
+      fprintf(mf, ",\n  \"unreplayable_apis\": [");
+      bool first = true;
+      for (const auto& api : g_unreplayable_apis) {
+        fprintf(mf, "%s\"%s\"", first ? "" : ", ", api.c_str());
+        first = false;
+      }
+      fprintf(mf, "]");
+    }
+  }
   if (g_metadata_json_len > 0) {
     fprintf(mf, ",\n  \"metadata\": %.*s\n",
             static_cast<int>(g_metadata_json_len), g_metadata_json);
@@ -737,19 +756,36 @@ void checkpoint() {
 void mark_incomplete(const char* reason) {
   // Record once; the loud, AMD_LOG_LEVEL-routed message is emitted by the caller
   // (e.g. serialize_kernel_launch) which has the relevant context. Here we only
-  // need the durable flag and a single stderr breadcrumb so a bare run still
-  // surfaces it.
+  // need the durable flag and a single breadcrumb so a bare run still surfaces
+  // it. Error level: the archive is not faithful and replay must not treat it
+  // as one. log_printf appends its own newline, so the format string omits it.
   if (!g_capture_incomplete.exchange(true, std::memory_order_relaxed)) {
-    fprintf(stderr,
-            "[HRR capture] Archive marked INCOMPLETE: %s. The clean-shutdown "
-            "trailer will be omitted and manifest.complete=false so replay "
-            "cannot treat this capture as faithful.\n",
-            reason ? reason : "(unspecified)");
+    LogPrintfError(
+        "[HRR capture] Archive marked INCOMPLETE: %s. The clean-shutdown "
+        "trailer will be omitted and manifest.complete=false so replay "
+        "cannot treat this capture as faithful",
+        reason ? reason : "(unspecified)");
   }
 }
 
 bool is_incomplete() { return g_capture_incomplete.load(std::memory_order_relaxed); }
 
+void note_unreplayable(const char* api, const char* reason) {
+  if (!api) return;
+  {
+    std::lock_guard<std::mutex> lk(g_unreplayable_mu);
+    if (!g_unreplayable_apis.insert(api).second) return;
+  }
+  // Warning, not Error: unlike mark_incomplete() the archive is well-formed and
+  // every event is present — only the ability to re-execute this one call is
+  // lost. That is a degradation, not a failure.
+  // log_printf appends its own newline, so the format string omits it.
+  LogPrintfWarning(
+      "[HRR capture] %s cannot be replayed: %s. The call is recorded, but "
+      "replay will report it as unreplayable rather than reproduce it",
+      api, reason ? reason : "(unspecified)");
+}
+
 void flush(const char* /*output_dir*/) {
   // Always finalize the *effective* directory this process actually wrote to
   // (g_output_dir), which is always a pid-<pid> sub-archive. The caller passes

projects/clr/hipamd/src/hrr/hip_capture_writer.h

@@ -71,6 +71,20 @@ void mark_incomplete(const char* reason);
 // Returns true if mark_incomplete() has been called.
 bool is_incomplete();
 
+// Record that `api` was captured but cannot be replayed faithfully — the call
+// is in the archive, its effect is not reproducible (a host callback belonging
+// to this process, a handle only meaningful to another rank, ...). Warns once
+// per API at capture time, so the gap is visible when the recording is made
+// rather than only when someone tries to replay it, and lists the API under
+// "unreplayable_apis" in manifest.json.
+//
+// Deliberately NOT mark_incomplete(): that means an event was dropped or torn
+// and suppresses the clean-shutdown trailer. Here every event is present and
+// the archive is well-formed; what is missing is the ability to re-execute
+// this particular call. Thread-safe; the warning fires once per API per
+// process.
+void note_unreplayable(const char* api, const char* reason);
+
 // Write a buffer as a content-addressed blob. Returns hash.
 // Thread-safe. Skips write if blob already exists on disk.
 Hash128 write_blob(const void* data, size_t len);