feat(amd): add synchronization tracing and correct capture data

example/amd/CMakeLists.txt

@@ -50,7 +50,8 @@ set(AMD_EXAMPLE_SOURCES
   gpufl_scope_demo.cpp
   memory_allocation_rows.cpp
   pm_sampling_sample_rows.cpp
-  vector_add_benchmark.cpp
+  synchronization_rows.cpp
+  vector_add_demo.cpp
 )
 
 set_source_files_properties(
@@ -62,18 +63,21 @@ hip_add_executable(amd_check_device check_device.cpp)
 hip_add_executable(amd_gpufl_scope_demo gpufl_scope_demo.cpp)
 hip_add_executable(amd_memory_allocation_rows memory_allocation_rows.cpp)
 hip_add_executable(amd_pm_sampling_sample_rows pm_sampling_sample_rows.cpp)
-hip_add_executable(amd_vector_add_benchmark vector_add_benchmark.cpp)
+hip_add_executable(amd_synchronization_rows synchronization_rows.cpp)
+hip_add_executable(amd_vector_add_demo vector_add_demo.cpp)
 
 if(TARGET hip::host)
   target_link_libraries(amd_check_device PRIVATE hip::host)
   target_link_libraries(amd_gpufl_scope_demo PRIVATE hip::host)
   target_link_libraries(amd_memory_allocation_rows PRIVATE hip::host)
   target_link_libraries(amd_pm_sampling_sample_rows PRIVATE hip::host)
-  target_link_libraries(amd_vector_add_benchmark PRIVATE hip::host)
+  target_link_libraries(amd_synchronization_rows PRIVATE hip::host)
+  target_link_libraries(amd_vector_add_demo PRIVATE hip::host)
 endif()
 
 target_link_libraries(amd_check_device PRIVATE gpufl::gpufl)
 target_link_libraries(amd_gpufl_scope_demo PRIVATE gpufl::gpufl)
 target_link_libraries(amd_memory_allocation_rows PRIVATE gpufl::gpufl)
 target_link_libraries(amd_pm_sampling_sample_rows PRIVATE gpufl::gpufl)
-target_link_libraries(amd_vector_add_benchmark PRIVATE gpufl::gpufl)
+target_link_libraries(amd_synchronization_rows PRIVATE gpufl::gpufl)
+target_link_libraries(amd_vector_add_demo PRIVATE gpufl::gpufl)

example/amd/README.md

@@ -9,6 +9,7 @@ This folder mirrors the CUDA example area with runnable HIP examples for AMD GPU
 - AMD kernel dispatch tracing via `rocprofiler-sdk`
 - AMD memcpy tracing via `rocprofiler-sdk`
 - AMD memory-allocation tracing via `rocprofiler-sdk`
+- AMD synchronization tracing via filtered ROCprofiler HIP runtime API records
 - Per-dispatch AMD hardware counters via ROCprofiler dispatch counting
 - Device-wide `PmSampling` timelines via ROCprofiler device counting
 - `gpufl` initialization with `backend = gpufl::BackendKind::Amd`
@@ -24,7 +25,7 @@ Today, the AMD backend is useful for:
 
 - system metric logging
 - device inventory
-- automatic HIP kernel, memcpy, and memory-allocation tracing
+- automatic HIP kernel, memcpy, memory-allocation, and synchronization tracing
 - per-dispatch and device-wide hardware-counter profiling
 - scope-level application instrumentation
 
@@ -36,14 +37,16 @@ It is not yet useful for:
 
 - `amd_check_device`
   - Basic HIP device detection smoke test
-- `amd_vector_add_benchmark`
-  - HIP vector add benchmark with result verification
+- `amd_vector_add_demo`
+  - Profiled HIP vector addition with one scoped kernel, H2D/D2H transfers, and result verification
 - `amd_gpufl_scope_demo`
   - Initializes `gpufl` with the AMD backend, runs HIP work inside scopes, and writes logs
 - `amd_memory_allocation_rows`
   - Exits successfully only when two HIP allocation/free phases emit memory-allocation rows
 - `amd_pm_sampling_sample_rows`
   - Selects AMD device counting and exits successfully only when each of two named scopes emits PM sample rows
+- `amd_synchronization_rows`
+  - Exits successfully only when two HIP workloads emit the expected synchronization rows
 
 ## Build
 
@@ -58,10 +61,11 @@ cmake -S . -B build-rocm-examples \
   -DBUILD_TESTING=OFF
 
 cmake --build build-rocm-examples --target amd_check_device
-cmake --build build-rocm-examples --target amd_vector_add_benchmark
+cmake --build build-rocm-examples --target amd_vector_add_demo
 cmake --build build-rocm-examples --target amd_gpufl_scope_demo
 cmake --build build-rocm-examples --target amd_memory_allocation_rows
 cmake --build build-rocm-examples --target amd_pm_sampling_sample_rows
+cmake --build build-rocm-examples --target amd_synchronization_rows
 ```
 
 The AMD example targets are only added when CMake detects HIP successfully.
@@ -115,17 +119,42 @@ subproject and disables the parent example/test targets to avoid recursion.
 
 ```bash
 ./build-rocm-examples/example/amd/amd_check_device
-./build-rocm-examples/example/amd/amd_vector_add_benchmark
+./build-rocm-examples/example/amd/amd_vector_add_demo
 ./build-rocm-examples/example/amd/amd_gpufl_scope_demo
 ./build-rocm-examples/example/amd/amd_memory_allocation_rows
 ./build-rocm-examples/example/amd/amd_pm_sampling_sample_rows
+./build-rocm-examples/example/amd/amd_synchronization_rows
 ```
 
+`amd_vector_add_demo` replaces the CPU-versus-GPU benchmark. It initializes
+GPUFlight in Trace mode, copies two 4 MiB input vectors to the GPU, launches
+`vectorAdd` inside `vector-addition-scope`, and copies the result back. It
+checks every output element and prints a session report; there are no CPU
+timings or speedup comparisons. Allocations and synchronization are also
+traced. A short capture may not contain periodic system metric samples.
+
+To include the demo source when running outside the repository (for example,
+from an IDE build directory), set the approved source root explicitly:
+
+```bash
+GPUFL_SOURCE_ROOT="$PWD/example/amd" \
+  ./build-rocm-examples/example/amd/amd_vector_add_demo
+```
+
+Run that command from the repository root. Source capture remains restricted
+to the approved directory; unrelated files and system headers are not included.
+
 `amd_memory_allocation_rows` enables `enable_memory_tracking`, runs two
 allocation/free phases, and checks that every expected allocate and free
 operation produces a `memory_alloc_event_batch` row. It returns exit code 2
 when ROCprofiler allocation tracing is unavailable or rows are missing.
 
+`amd_synchronization_rows` issues event synchronize, stream wait-event,
+stream synchronize, and device synchronize calls in each phase, then checks
+that the synchronization row count increases by the expected amount. It
+returns exit code 2 when ROCprofiler HIP runtime tracing is unavailable or
+rows are missing.
+
 `amd_pm_sampling_sample_rows` requests the portable `GPUBusy` counter, runs
 GPU work in `pm_rows_phase_a` and `pm_rows_phase_b`, and checks that the PM row
 count increases after each scope. It returns exit code 2 when AMD device
@@ -153,6 +182,12 @@ Success! Device 0: AMD Radeon RX 9070 XT (arch gfx1201, capability 12.0)
 
 ## Logs
 
+`amd_vector_add_demo` writes logs with prefix:
+
+```text
+gfl_amd_vector_add
+```
+
 `amd_gpufl_scope_demo` writes logs with prefix:
 
 ```bash
@@ -171,6 +206,25 @@ gfl_amd_memory_rows
 gfl_amd_pm_rows
 ```
 
+`amd_synchronization_rows` writes logs with prefix:
+
+```bash
+gfl_amd_sync_rows
+```
+
+Check a completed example capture before uploading it:
+
+```bash
+python3 example/amd/verify_trace_timestamps.py gfl_amd_pm_rows/<session-id>
+```
+
+This checks that kernel, copy, allocation, synchronization, scope, and PM rows
+use the session epoch clock. Use the corresponding log directory for other
+examples. Static ISA mappings are checked for duplicate delivery, not timing.
+Run `python3 example/amd/test_verify_trace_timestamps.py` for the validator tests. Old
+captures recorded with profiler-relative timestamps must be regenerated;
+re-uploading those same files will not repair their clock.
+
 With `rocprofiler-sdk` available, expect:
 
 - `job_start` inventory
@@ -179,6 +233,7 @@ With `rocprofiler-sdk` available, expect:
 - `kernel_detail`
 - `memcpy_event_batch`
 - `memory_alloc_event_batch` when `enable_memory_tracking` is enabled
+- `synchronization_event_batch` when `enable_synchronization` is enabled
 - `profile_sample_batch` for dispatch-counting requests
 - `pm_sampling_config` and `pm_sample_batch` for `PmSampling`
 - system metric samples

example/amd/synchronization_rows.cpp

@@ -0,0 +1,215 @@
+#include <hip/hip_runtime.h>
+
+#include <chrono>
+#include <cstdint>
+#include <iostream>
+#include <thread>
+
+#include "gpufl/core/monitor.hpp"
+#include "gpufl/gpufl.hpp"
+
+namespace {
+
+constexpr int kElementCount = 1 << 18;
+constexpr int kBlockSize = 256;
+constexpr int kIterations = 256;
+constexpr uint64_t kExpectedRowsPerPhase = 4;
+constexpr auto kDeliveryTimeout = std::chrono::seconds(2);
+
+bool CheckHip(const hipError_t status, const char* what) {
+    if (status == hipSuccess) return true;
+    std::cerr << what << " failed: " << hipGetErrorString(status) << "\n";
+    return false;
+}
+
+__global__ void synchronizationWorkload(float* values, const int count,
+                                        const int iterations) {
+    const int index = blockIdx.x * blockDim.x + threadIdx.x;
+    if (index >= count) return;
+
+    float value = values[index] + static_cast<float>((index & 255) + 1);
+    for (int iteration = 0; iteration < iterations; ++iteration) {
+        value = value * 1.000001f + 0.000001f;
+        if (value > 4096.0f) value -= 4096.0f;
+    }
+    values[index] = value;
+}
+
+bool RunSynchronizationPhase(float* values, const hipStream_t producer,
+                             const hipStream_t consumer,
+                             const hipEvent_t event) {
+    const dim3 block(kBlockSize);
+    const dim3 grid((kElementCount + block.x - 1) / block.x);
+
+    hipLaunchKernelGGL(synchronizationWorkload, grid, block, 0, producer,
+                       values, kElementCount, kIterations);
+    if (!CheckHip(hipGetLastError(), "producer workload launch")) return false;
+    if (!CheckHip(hipEventRecord(event, producer), "hipEventRecord")) {
+        return false;
+    }
+    if (!CheckHip(hipStreamWaitEvent(consumer, event, 0),
+                  "hipStreamWaitEvent")) {
+        return false;
+    }
+
+    hipLaunchKernelGGL(synchronizationWorkload, grid, block, 0, consumer,
+                       values, kElementCount, kIterations);
+    if (!CheckHip(hipGetLastError(), "consumer workload launch")) return false;
+
+    if (!CheckHip(hipStreamSynchronize(consumer), "hipStreamSynchronize")) {
+        return false;
+    }
+    if (!CheckHip(hipEventSynchronize(event), "hipEventSynchronize")) {
+        return false;
+    }
+    return CheckHip(hipDeviceSynchronize(), "hipDeviceSynchronize");
+}
+
+uint64_t WaitForSynchronizationRows(const uint64_t minimum_rows) {
+    const auto deadline = std::chrono::steady_clock::now() + kDeliveryTimeout;
+    uint64_t rows = gpufl::Monitor::SynchronizationRowsSeen();
+    while (rows < minimum_rows && std::chrono::steady_clock::now() < deadline) {
+        std::this_thread::sleep_for(std::chrono::milliseconds(10));
+        rows = gpufl::Monitor::SynchronizationRowsSeen();
+    }
+    return rows;
+}
+
+}  // namespace
+
+int main() {
+    gpufl::InitOptions opts;
+    opts.app_name = "amd_synchronization_rows";
+    opts.log_path = "gfl_amd_sync_rows";
+    opts.backend = gpufl::BackendKind::Amd;
+    opts.profiling_engine = gpufl::ProfilingEngine::Trace;
+    opts.enable_synchronization = true;
+    opts.enable_memory_tracking = false;
+    opts.continuous_system_sampling = false;
+    opts.enable_debug_output = true;
+    opts.enable_stack_trace = false;
+
+    if (!gpufl::init(opts)) {
+        std::cerr << "Failed to initialize gpufl for AMD synchronization tracing\n";
+        return 1;
+    }
+
+    const std::string engine =
+        gpufl::Monitor::ResolvedProfilingEngineWireName();
+    const bool engine_ok = engine == "amd.buffer_tracing";
+    std::cout << "=== GPUFL AMD Synchronization Rows ===\n"
+              << "Resolved engine: " << engine << "\n";
+    if (!engine_ok) {
+        std::cerr << "Expected amd.buffer_tracing; ROCprofiler tracing may be unavailable\n";
+    }
+
+    float* device_values = nullptr;
+    hipStream_t producer = nullptr;
+    hipStream_t consumer = nullptr;
+    hipEvent_t event = nullptr;
+
+    bool workload_ok = CheckHip(
+        hipMalloc(&device_values,
+                  static_cast<size_t>(kElementCount) * sizeof(float)),
+        "hipMalloc(device_values)");
+    if (workload_ok) {
+        workload_ok = CheckHip(
+            hipMemset(device_values, 0,
+                      static_cast<size_t>(kElementCount) * sizeof(float)),
+            "hipMemset(device_values)");
+    }
+    if (workload_ok) {
+        workload_ok = CheckHip(
+            hipStreamCreateWithFlags(&producer, hipStreamNonBlocking),
+            "hipStreamCreate(producer)");
+    }
+    if (workload_ok) {
+        workload_ok = CheckHip(
+            hipStreamCreateWithFlags(&consumer, hipStreamNonBlocking),
+            "hipStreamCreate(consumer)");
+    }
+    if (workload_ok) {
+        workload_ok = CheckHip(
+            hipEventCreateWithFlags(&event, hipEventDisableTiming),
+            "hipEventCreate");
+    }
+
+    // The first HIP calls load HSA. GPUFlight then retries its deferred
+    // ROCprofiler context start from the collector thread.
+    std::this_thread::sleep_for(std::chrono::milliseconds(50));
+
+    const uint64_t initial_rows =
+        gpufl::Monitor::SynchronizationRowsSeen();
+    bool priming_ok = false;
+    uint64_t rows_before = initial_rows;
+    if (workload_ok) {
+        priming_ok = RunSynchronizationPhase(device_values, producer,
+                                             consumer, event);
+        if (priming_ok) {
+            rows_before = WaitForSynchronizationRows(
+                initial_rows + kExpectedRowsPerPhase);
+            std::this_thread::sleep_for(std::chrono::milliseconds(20));
+            rows_before = gpufl::Monitor::SynchronizationRowsSeen();
+        }
+    }
+    const bool priming_rows =
+        rows_before >= initial_rows + kExpectedRowsPerPhase;
+    if (!priming_rows) {
+        std::cerr << "Priming phase did not emit all synchronization rows\n";
+    }
+    workload_ok = workload_ok && priming_ok && priming_rows;
+
+    bool phase_a_ok = false;
+    if (workload_ok) {
+        GFL_SCOPE("sync_rows_phase_a") {
+            phase_a_ok = RunSynchronizationPhase(device_values, producer,
+                                                 consumer, event);
+        }
+    }
+    const uint64_t rows_after_a = WaitForSynchronizationRows(
+        rows_before + kExpectedRowsPerPhase);
+
+    bool phase_b_ok = false;
+    if (workload_ok && phase_a_ok) {
+        GFL_SCOPE("sync_rows_phase_b") {
+            phase_b_ok = RunSynchronizationPhase(device_values, producer,
+                                                 consumer, event);
+        }
+    }
+    const uint64_t rows_after_b = WaitForSynchronizationRows(
+        rows_after_a + kExpectedRowsPerPhase);
+
+    std::cout << "Synchronization rows: before=" << rows_before
+              << ", after phase A=" << rows_after_a
+              << ", after phase B=" << rows_after_b << "\n";
+
+    const bool phase_a_rows =
+        rows_after_a >= rows_before + kExpectedRowsPerPhase;
+    const bool phase_b_rows =
+        rows_after_b >= rows_after_a + kExpectedRowsPerPhase;
+    if (!phase_a_rows) {
+        std::cerr << "Phase A did not emit all synchronization rows\n";
+    }
+    if (!phase_b_rows) {
+        std::cerr << "Phase B did not emit all synchronization rows\n";
+    }
+
+    if (event != nullptr) (void)hipEventDestroy(event);
+    if (consumer != nullptr) (void)hipStreamDestroy(consumer);
+    if (producer != nullptr) (void)hipStreamDestroy(producer);
+    if (device_values != nullptr) (void)hipFree(device_values);
+
+    gpufl::shutdown();
+    gpufl::generateReport();
+
+    const bool passed = engine_ok && workload_ok && priming_ok &&
+                        priming_rows && phase_a_ok && phase_b_ok &&
+                        phase_a_rows && phase_b_rows;
+    if (!passed) return 2;
+
+    std::cout
+        << "\nPASS: both phases emitted the four expected AMD synchronization rows.\n"
+        << "Inspect logs with prefix " << opts.log_path
+        << " for synchronization_event_batch events.\n";
+    return 0;
+}