From 35059818368efbb59f027ade656edb6e4ff7fe98 Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Wed, 9 Sep 2026 15:05:54 -0700 Subject: [PATCH] feat(amd): add synchronization tracing and correct capture data --- example/amd/CMakeLists.txt | 12 +- example/amd/README.md | 65 +++++- example/amd/synchronization_rows.cpp | 215 ++++++++++++++++++ example/amd/test_verify_trace_timestamps.py | 36 +++ example/amd/vector_add_benchmark.cpp | 137 ----------- example/amd/vector_add_demo.cpp | 98 ++++++++ example/amd/verify_trace_timestamps.py | 79 +++++++ .../backends/amd/amd_capture_capabilities.cpp | 23 ++ .../backends/amd/amd_capture_capabilities.hpp | 3 + .../gpufl/backends/amd/amd_trace_policy.cpp | 15 ++ .../gpufl/backends/amd/amd_trace_policy.hpp | 13 ++ .../backends/amd/rocprofiler_backend.cpp | 203 ++++++++++++++++- .../backends/amd/rocprofiler_backend.hpp | 15 ++ include/gpufl/core/dictionary_manager.cpp | 3 +- include/gpufl/core/monitor.cpp | 4 + include/gpufl/core/monitor.hpp | 10 +- include/gpufl/core/monitor_batch_manager.cpp | 6 + include/gpufl/core/monitor_batch_manager.hpp | 3 + include/gpufl/gpufl.hpp | 20 +- include/gpufl/report/text_report.cpp | 85 +++++++ include/gpufl/report/text_report.hpp | 12 + .../amd/test_amd_profiling_policy.cpp | 38 ++++ tests/backends/amd/test_amd_trace_policy.cpp | 16 ++ tests/core/test_monitor.cpp | 18 ++ tests/core/test_text_report.cpp | 38 ++++ 25 files changed, 1000 insertions(+), 167 deletions(-) create mode 100644 example/amd/synchronization_rows.cpp create mode 100644 example/amd/test_verify_trace_timestamps.py delete mode 100644 example/amd/vector_add_benchmark.cpp create mode 100644 example/amd/vector_add_demo.cpp create mode 100644 example/amd/verify_trace_timestamps.py diff --git a/example/amd/CMakeLists.txt b/example/amd/CMakeLists.txt index 91b5dfe..b7028e7 100644 --- a/example/amd/CMakeLists.txt +++ b/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) diff --git a/example/amd/README.md b/example/amd/README.md index 0b0878b..81f961b 100644 --- a/example/amd/README.md +++ b/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/ +``` + +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 diff --git a/example/amd/synchronization_rows.cpp b/example/amd/synchronization_rows.cpp new file mode 100644 index 0000000..862b323 --- /dev/null +++ b/example/amd/synchronization_rows.cpp @@ -0,0 +1,215 @@ +#include + +#include +#include +#include +#include + +#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((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(kElementCount) * sizeof(float)), + "hipMalloc(device_values)"); + if (workload_ok) { + workload_ok = CheckHip( + hipMemset(device_values, 0, + static_cast(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; +} diff --git a/example/amd/test_verify_trace_timestamps.py b/example/amd/test_verify_trace_timestamps.py new file mode 100644 index 0000000..eeab54b --- /dev/null +++ b/example/amd/test_verify_trace_timestamps.py @@ -0,0 +1,36 @@ +"""Portable regression tests for the example capture validator.""" +import unittest +from verify_trace_timestamps import validate + + +class CaptureValidationTest(unittest.TestCase): + def events(self): + epoch = 1788928712000000000 + return [ + {"type": "job_start", "session_id": "test", "ts_ns": epoch}, + {"type": "shutdown", "ts_ns": epoch + 10000000}, + {"type": "kernel_event_batch", "base_time_ns": epoch, + "columns": ["dt_ns", "duration_ns"], "rows": [[1000, 100]]}, + ] + + def test_rejects_profiler_relative_timestamps(self): + for kind in ("kernel_event_batch", "memcpy_event_batch", + "memory_alloc_event_batch", "synchronization_event_batch"): + with self.subTest(kind=kind): + events = self.events() + events[-1]["type"] = kind + validate(events) + events[-1]["base_time_ns"] = 491751431486779 + with self.assertRaisesRegex(ValueError, "outside the session clock"): + validate(events) + + def test_static_batch_must_not_be_broadcast_to_every_channel(self): + batch = {"type": "profile_sample_batch", "session_id": "test", "batch_id": 1000001, + "columns": ["sample_kind"], "rows": [[2]]} + validate(self.events() + [batch]) + with self.assertRaisesRegex(ValueError, "Duplicate static ISA"): + validate(self.events() + [batch, batch, batch, batch]) + + +if __name__ == "__main__": + unittest.main() diff --git a/example/amd/vector_add_benchmark.cpp b/example/amd/vector_add_benchmark.cpp deleted file mode 100644 index 21f715d..0000000 --- a/example/amd/vector_add_benchmark.cpp +++ /dev/null @@ -1,137 +0,0 @@ -#include -#include -#include - -#include - -static 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 vectorAddGPU(const int* a, const int* b, int* c, int n) { - int i = blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) { - c[i] = a[i] + b[i]; - } -} - -void vectorAddCPU(const int* a, const int* b, int* c, int n) { - for (int i = 0; i < n; i++) { - c[i] = a[i] + b[i]; - } -} - -int main() { - const int n = 1 << 24; - const size_t size = static_cast(n) * sizeof(int); - - std::cout << "HIP Vector Addition Benchmark (n = " << n << " elements)\n"; - std::cout << "Data size: " << static_cast(size) / (1024 * 1024) - << " MB per vector\n"; - - int* h_a = static_cast(std::malloc(size)); - int* h_b = static_cast(std::malloc(size)); - int* h_c_cpu = static_cast(std::malloc(size)); - int* h_c_gpu = static_cast(std::malloc(size)); - - if (!h_a || !h_b || !h_c_cpu || !h_c_gpu) { - std::cerr << "Failed to allocate host memory\n"; - return 1; - } - - for (int i = 0; i < n; i++) { - h_a[i] = i; - h_b[i] = i * 2; - } - - auto start_cpu = std::chrono::high_resolution_clock::now(); - vectorAddCPU(h_a, h_b, h_c_cpu, n); - auto end_cpu = std::chrono::high_resolution_clock::now(); - std::chrono::duration cpu_time = end_cpu - start_cpu; - std::cout << "CPU time: " << cpu_time.count() << " ms\n"; - - int* d_a = nullptr; - int* d_b = nullptr; - int* d_c = nullptr; - if (!checkHip(hipMalloc(&d_a, size), "hipMalloc(d_a)") || - !checkHip(hipMalloc(&d_b, size), "hipMalloc(d_b)") || - !checkHip(hipMalloc(&d_c, size), "hipMalloc(d_c)")) { - return 1; - } - - hipEvent_t start_event{}; - hipEvent_t stop_event{}; - if (!checkHip(hipEventCreate(&start_event), "hipEventCreate(start)") || - !checkHip(hipEventCreate(&stop_event), "hipEventCreate(stop)")) { - return 1; - } - - auto start_total = std::chrono::high_resolution_clock::now(); - - if (!checkHip(hipMemcpy(d_a, h_a, size, hipMemcpyHostToDevice), - "hipMemcpy H2D a") || - !checkHip(hipMemcpy(d_b, h_b, size, hipMemcpyHostToDevice), - "hipMemcpy H2D b")) { - return 1; - } - - const int threadsPerBlock = 256; - const int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock; - - if (!checkHip(hipEventRecord(start_event), "hipEventRecord(start)")) { - return 1; - } - hipLaunchKernelGGL(vectorAddGPU, dim3(blocksPerGrid), dim3(threadsPerBlock), - 0, 0, d_a, d_b, d_c, n); - if (!checkHip(hipGetLastError(), "vectorAddGPU launch") || - !checkHip(hipEventRecord(stop_event), "hipEventRecord(stop)")) { - return 1; - } - - if (!checkHip(hipMemcpy(h_c_gpu, d_c, size, hipMemcpyDeviceToHost), - "hipMemcpy D2H c")) { - return 1; - } - - auto end_total = std::chrono::high_resolution_clock::now(); - std::chrono::duration gpu_total_time = - end_total - start_total; - - float kernel_time = 0.0f; - if (!checkHip(hipEventSynchronize(stop_event), "hipEventSynchronize(stop)") || - !checkHip(hipEventElapsedTime(&kernel_time, start_event, stop_event), - "hipEventElapsedTime")) { - return 1; - } - - std::cout << "GPU Kernel time: " << kernel_time << " ms\n"; - std::cout << "GPU Total time (including H2D and D2H): " - << gpu_total_time.count() << " ms\n"; - std::cout << "Speedup (CPU / GPU Kernel): " - << cpu_time.count() / kernel_time << "x\n"; - std::cout << "Speedup (CPU / GPU Total): " - << cpu_time.count() / gpu_total_time.count() << "x\n"; - - bool passed = true; - for (int i = 0; i < n; i++) { - if (h_c_cpu[i] != h_c_gpu[i]) { - passed = false; - break; - } - } - std::cout << "Verification: " << (passed ? "PASSED" : "FAILED") << "\n"; - - (void)hipFree(d_a); - (void)hipFree(d_b); - (void)hipFree(d_c); - (void)hipEventDestroy(start_event); - (void)hipEventDestroy(stop_event); - std::free(h_a); - std::free(h_b); - std::free(h_c_cpu); - std::free(h_c_gpu); - - return passed ? 0 : 1; -} diff --git a/example/amd/vector_add_demo.cpp b/example/amd/vector_add_demo.cpp new file mode 100644 index 0000000..9f40809 --- /dev/null +++ b/example/amd/vector_add_demo.cpp @@ -0,0 +1,98 @@ +#include + +#include +#include + +#include "gpufl/gpufl.hpp" + +namespace { + +constexpr int kElementCount = 1 << 20; +constexpr int kBlockSize = 256; + +bool CheckHip(const hipError_t status, const char* what) { + if (status == hipSuccess) return true; + std::cerr << what << " failed: " << hipGetErrorString(status) << "\n"; + return false; +} + +} // namespace + +__global__ void vectorAdd(const float* a, const float* b, float* c, int count) { + const int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < count) c[i] = a[i] + b[i]; +} + +int main() { + const size_t bytes = static_cast(kElementCount) * sizeof(float); + std::vector a(kElementCount), b(kElementCount), c(kElementCount); + for (int i = 0; i < kElementCount; ++i) { + // Small integer-valued floats make exact result verification possible. + a[i] = static_cast(i % 1024); + b[i] = static_cast((i % 256) * 2); + } + + gpufl::InitOptions opts; + opts.app_name = "amd_vector_add_demo"; + opts.log_path = "gfl_amd_vector_add"; + opts.backend = gpufl::BackendKind::Amd; + opts.profiling_engine = gpufl::ProfilingEngine::Trace; + opts.enable_memory_tracking = true; + opts.enable_synchronization = true; + opts.continuous_system_sampling = true; + opts.system_sample_rate_ms = 50; + opts.enable_stack_trace = false; + if (!gpufl::init(opts)) { + std::cerr << "Failed to initialize GPUFlight for AMD vector addition\n"; + return 1; + } + + float* d_a = nullptr; + float* d_b = nullptr; + float* d_c = nullptr; + bool ok = CheckHip(hipMalloc(&d_a, bytes), "hipMalloc(a)") && + CheckHip(hipMalloc(&d_b, bytes), "hipMalloc(b)") && + CheckHip(hipMalloc(&d_c, bytes), "hipMalloc(c)"); + if (ok) { + ok = CheckHip(hipMemcpy(d_a, a.data(), bytes, hipMemcpyHostToDevice), + "H2D a") && + CheckHip(hipMemcpy(d_b, b.data(), bytes, hipMemcpyHostToDevice), + "H2D b"); + } + if (ok) { + GFL_SCOPE("vector-addition-scope") { + const dim3 block(kBlockSize); + const dim3 grid((kElementCount + kBlockSize - 1) / kBlockSize); + hipLaunchKernelGGL(vectorAdd, grid, block, 0, 0, + d_a, d_b, d_c, kElementCount); + ok = CheckHip(hipGetLastError(), "vectorAdd launch") && + CheckHip(hipDeviceSynchronize(), "vectorAdd synchronize"); + } + } + if (ok) { + ok = CheckHip(hipMemcpy(c.data(), d_c, bytes, hipMemcpyDeviceToHost), + "D2H c"); + } + if (ok) { + for (int i = 0; i < kElementCount; ++i) { + if (c[i] != a[i] + b[i]) { + std::cerr << "Verification failed at element " << i << "\n"; + ok = false; + break; + } + } + } + + // Always release successful allocations, including after a partial failure. + // Keep the profiler running until frees have been captured. + if (d_c) ok = CheckHip(hipFree(d_c), "hipFree(c)") && ok; + if (d_b) ok = CheckHip(hipFree(d_b), "hipFree(b)") && ok; + if (d_a) ok = CheckHip(hipFree(d_a), "hipFree(a)") && ok; + gpufl::shutdown(); + gpufl::generateReport(); + if (!ok) return 2; + + std::cout << "\nPASS: vectorAdd verified all " << kElementCount + << " elements.\nLogs: " << opts.log_path << "\n"; + return 0; +} diff --git a/example/amd/verify_trace_timestamps.py b/example/amd/verify_trace_timestamps.py new file mode 100644 index 0000000..5edc3a3 --- /dev/null +++ b/example/amd/verify_trace_timestamps.py @@ -0,0 +1,79 @@ +#!/usr/bin/env python3 +"""Check one completed AMD example session's timeline clock without uploading it. + +Usage: python3 example/amd/verify_trace_timestamps.py +Static ISA mappings are checked for duplicate delivery, not for timestamps. +""" +import argparse +from collections import Counter +import gzip +import json +from pathlib import Path + + +TIMED_BATCHES = { + "kernel_event_batch", "memcpy_event_batch", "memory_alloc_event_batch", + "synchronization_event_batch", "scope_event_batch", "pm_sample_batch", +} + + +def validate(events): + session_ids = {e["session_id"] for e in events if "session_id" in e} + if len(session_ids) != 1: + raise ValueError("Provide the log directory for exactly one session") + starts = [e["ts_ns"] for e in events if e["type"] == "job_start"] + ends = [e["ts_ns"] for e in events if e["type"] == "shutdown"] + if not starts or not ends: + raise ValueError("Expected job_start and shutdown in a completed capture") + # Tracing can start before job_start is emitted, including HIP's internal + # allocations. A one-second margin admits setup, not a different epoch. + lower, upper = min(starts) - 1_000_000_000, max(ends) + 1_000_000_000 + static_batches = set() + counts = Counter() + timestamps = [] + for event in events: + kind = event["type"] + if kind == "profile_sample_batch": + rows = [dict(zip(event["columns"], values)) for values in event["rows"]] + if any(row.get("sample_kind") == 2 for row in rows): + key = (event["session_id"], event["batch_id"]) + if key in static_batches: + raise ValueError("Duplicate static ISA batch across log channels") + static_batches.add(key) + if kind not in TIMED_BATCHES: + continue + for values in event["rows"]: + row = dict(zip(event["columns"], values)) + start = event["base_time_ns"] + row["dt_ns"] + duration = row.get("duration_ns", 0) + if duration < 0 or not lower <= start <= start + duration <= upper: + raise ValueError(f"{kind}: timestamp {start} / duration {duration} outside the session clock") + counts[kind] += 1 + timestamps.extend((start, start + duration)) + if not counts: + raise ValueError("No timed activity rows found") + return counts, (max(timestamps) - min(timestamps)) / 1_000_000 + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("session_dir", type=Path) + args = parser.parse_args() + events = [] + for path in sorted(args.session_dir.iterdir()): + if path.name.endswith(".log.gz"): + opener = gzip.open + elif path.name.endswith(".log"): + opener = open + else: + continue + with opener(path, "rt", encoding="utf-8") as stream: + events.extend(json.loads(line) for line in stream if line.strip()) + counts, span_ms = validate(events) + for kind, count in sorted(counts.items()): + print(f"{kind}: {count} rows") + print(f"PASS: all timed rows share the session epoch; activity span {span_ms:.3f} ms") + + +if __name__ == "__main__": + main() diff --git a/include/gpufl/backends/amd/amd_capture_capabilities.cpp b/include/gpufl/backends/amd/amd_capture_capabilities.cpp index 16bf857..51ad561 100644 --- a/include/gpufl/backends/amd/amd_capture_capabilities.cpp +++ b/include/gpufl/backends/amd/amd_capture_capabilities.cpp @@ -89,6 +89,29 @@ CaptureCapabilitiesEvent BuildAmdCaptureCapabilitiesEvent( ? "Memory-copy tracing was enabled but emitted no rows this session." : "Memory-copy tracing was not active.")); + AddCapability( + event, "sync_activity", input.synchronization_requested, + !input.synchronization_requested + ? "not_requested" + : input.synchronization_configured + ? (input.synchronization_rows > 0 ? "collected" : "enabled_no_data") + : "skipped", + input.synchronization_configured + ? "rocprofiler_hip_runtime_api_ext" + : "disabled", + input.synchronization_configured + ? (input.synchronization_rows > 0 ? "" : "enabled_but_no_records") + : input.synchronization_requested + ? "rocprofiler_hip_runtime_api_unavailable" + : "", + input.synchronization_configured + ? (input.synchronization_rows > 0 + ? "HIP synchronization records were collected through ROCprofiler SDK." + : "HIP synchronization tracing was enabled but emitted no rows this session.") + : input.synchronization_requested + ? "HIP synchronization tracing was requested but unavailable." + : "Synchronization tracing was not requested."); + AddCapability( event, "memory_activity", input.memory_activity_requested, !input.memory_activity_requested diff --git a/include/gpufl/backends/amd/amd_capture_capabilities.hpp b/include/gpufl/backends/amd/amd_capture_capabilities.hpp index a9ad61d..c2ab495 100644 --- a/include/gpufl/backends/amd/amd_capture_capabilities.hpp +++ b/include/gpufl/backends/amd/amd_capture_capabilities.hpp @@ -12,10 +12,13 @@ struct AmdCaptureCapabilityInput { int64_t ts_ns = 0; AmdResolvedProfilingPlan plan; bool trace_configured = false; + bool synchronization_requested = false; + bool synchronization_configured = false; bool memory_activity_requested = false; bool memory_activity_configured = false; uint64_t kernel_rows = 0; uint64_t memcpy_rows = 0; + uint64_t synchronization_rows = 0; uint64_t memory_activity_rows = 0; uint64_t profiling_sample_rows = 0; uint64_t pm_sample_rows = 0; diff --git a/include/gpufl/backends/amd/amd_trace_policy.cpp b/include/gpufl/backends/amd/amd_trace_policy.cpp index 11a8aa3..7ee5b0b 100644 --- a/include/gpufl/backends/amd/amd_trace_policy.cpp +++ b/include/gpufl/backends/amd/amd_trace_policy.cpp @@ -40,4 +40,19 @@ uint8_t ResolveAmdMemoryAllocationKind(const AmdTraceEndpoint& agent) { return agent.kind == AmdTraceAgentKind::Gpu ? uint8_t{3} : uint8_t{0}; } +uint8_t ResolveAmdSynchronizationType( + const AmdSynchronizationOperation operation) { + switch (operation) { + case AmdSynchronizationOperation::EventSynchronize: + return 1; + case AmdSynchronizationOperation::StreamWaitEvent: + return 2; + case AmdSynchronizationOperation::StreamSynchronize: + return 3; + case AmdSynchronizationOperation::ContextSynchronize: + return 4; + } + return 0; +} + } // namespace gpufl::amd diff --git a/include/gpufl/backends/amd/amd_trace_policy.hpp b/include/gpufl/backends/amd/amd_trace_policy.hpp index bc6a9a8..cb0859d 100644 --- a/include/gpufl/backends/amd/amd_trace_policy.hpp +++ b/include/gpufl/backends/amd/amd_trace_policy.hpp @@ -43,4 +43,17 @@ std::optional NormalizeAmdMemoryAllocationOperation( // classified as DEVICE (3) while CPU allocations remain UNKNOWN (0). uint8_t ResolveAmdMemoryAllocationKind(const AmdTraceEndpoint& agent); +// ROCprofiler identifies synchronization through HIP runtime API operations. +// Keep native operation ids out of the portable policy layer and map the +// classified operation to GPUFlight's CUPTI-compatible wire values here. +enum class AmdSynchronizationOperation { + EventSynchronize, + StreamWaitEvent, + StreamSynchronize, + ContextSynchronize, +}; + +uint8_t ResolveAmdSynchronizationType( + AmdSynchronizationOperation operation); + } // namespace gpufl::amd diff --git a/include/gpufl/backends/amd/rocprofiler_backend.cpp b/include/gpufl/backends/amd/rocprofiler_backend.cpp index e4563ce..199a37b 100644 --- a/include/gpufl/backends/amd/rocprofiler_backend.cpp +++ b/include/gpufl/backends/amd/rocprofiler_backend.cpp @@ -20,6 +20,7 @@ #include #include #include +#include #include #include @@ -86,6 +87,32 @@ const char* CopyKindName(const uint32_t kind) { } } +std::optional ClassifySynchronizationOperation( + const rocprofiler_tracing_operation_t operation) { + switch (operation) { + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipEventSynchronize): + return AmdSynchronizationOperation::EventSynchronize; + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent): + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent_spt): + return AmdSynchronizationOperation::StreamWaitEvent; + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamSynchronize): + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamSynchronize_spt): + return AmdSynchronizationOperation::StreamSynchronize; + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipCtxSynchronize): + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipDeviceSynchronize): + return AmdSynchronizationOperation::ContextSynchronize; + default: + return std::nullopt; + } +} + } // namespace bool RocprofilerBackend::IsAvailable(std::string* reason) { @@ -121,6 +148,7 @@ void RocprofilerBackend::initialize(const MonitorOptions& opts) { opts_ = opts; kernel_rows_emitted_.store(0, std::memory_order_relaxed); memcpy_rows_emitted_.store(0, std::memory_order_relaxed); + synchronization_rows_emitted_.store(0, std::memory_order_relaxed); memory_activity_rows_emitted_.store(0, std::memory_order_relaxed); trace_records_dropped_.store(0, std::memory_order_relaxed); trace_records_queue_dropped_.store(0, std::memory_order_relaxed); @@ -133,6 +161,7 @@ void RocprofilerBackend::initialize(const MonitorOptions& opts) { capture_capabilities_session_id_.clear(); capability_kernel_rows_baseline_ = 0; capability_memcpy_rows_baseline_ = 0; + capability_synchronization_rows_baseline_ = 0; capability_memory_activity_rows_baseline_ = 0; capability_pm_sample_rows_baseline_ = 0; capability_dropped_records_baseline_ = 0; @@ -158,6 +187,18 @@ bool RocprofilerBackend::configureRocprofiler(const MonitorOptions& opts, std::string* reason) { (void) opts; if (!IsAvailable(reason)) return false; + // ROCprofiler timestamps are not Unix timestamps. Bracket a clock sample + // with our epoch clock and use the midpoint to minimize calibration skew. + // Keep this mapping stable for the lifetime of this capture, including + // records delivered after a context stop. + rocprofiler_timestamp_t profiler_ns = 0; + const int64_t before_ns = detail::GetTimestampNs(); + const auto status = rocprofiler_get_timestamp(&profiler_ns); + const int64_t after_ns = detail::GetTimestampNs(); + if (!CheckStatus(status, "rocprofiler_get_timestamp", reason)) return false; + trace_epoch_offset_ns_ = before_ns + (after_ns - before_ns) / 2 + - static_cast(profiler_ns); + if (!registerTool(reason)) return false; return true; } @@ -168,6 +209,7 @@ void RocprofilerBackend::resetToolState() { client_handle_ = 0; client_finalize_ = nullptr; tool_registered_.store(false); + synchronization_configured_.store(false); memory_activity_configured_.store(false); active_.store(false); start_requested_.store(false); @@ -175,17 +217,24 @@ void RocprofilerBackend::resetToolState() { deferred_start_logged_.store(false); start_failure_logged_.store(false); { - std::lock_guard lock(kernel_meta_mutex_); + std::lock_guard lock(kernel_meta_mutex_); kernel_metadata_.clear(); } { - std::lock_guard lock(external_scope_mutex_); + std::lock_guard lock(external_scope_mutex_); external_scope_metadata_.clear(); } { - std::lock_guard lock(memory_allocation_mutex_); + std::lock_guard lock(memory_allocation_mutex_); memory_allocations_.clear(); } + { + std::lock_guard lock(synchronization_handle_mutex_); + hip_stream_ids_.clear(); + hip_event_ids_.clear(); + next_hip_stream_id_ = 1; + next_hip_event_id_ = 1; + } { std::lock_guard lock(agent_mutex_); gpu_device_ids_.clear(); @@ -317,6 +366,37 @@ int RocprofilerBackend::toolInitialize() { &reason)) { return -1; } + if (opts_.enable_synchronization) { + const std::array + synchronization_operations = { + static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipEventSynchronize), + static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent), + static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent_spt), + static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamSynchronize), + static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamSynchronize_spt), + static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipCtxSynchronize), + static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipDeviceSynchronize), + }; + const auto status = rocprofiler_configure_buffer_tracing_service( + context_, ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API_EXT, + synchronization_operations.data(), + synchronization_operations.size(), buffer_); + if (status == ROCPROFILER_STATUS_SUCCESS) { + synchronization_configured_.store(true, + std::memory_order_release); + } else { + GFL_LOG_WARN( + "[ROCProfilerBackend] HIP synchronization tracing unavailable: ", + StatusToString(status)); + } + } if (opts_.enable_memory_tracking) { const auto status = rocprofiler_configure_buffer_tracing_service( @@ -538,6 +618,8 @@ void RocprofilerBackend::emitCapabilities() { kernel_rows_emitted_.load(std::memory_order_relaxed); const uint64_t memcpy_rows = memcpy_rows_emitted_.load(std::memory_order_relaxed); + const uint64_t synchronization_rows = + synchronization_rows_emitted_.load(std::memory_order_relaxed); const uint64_t memory_activity_rows = memory_activity_rows_emitted_.load(std::memory_order_relaxed); const uint64_t dropped_records = @@ -561,10 +643,15 @@ void RocprofilerBackend::emitCapabilities() { input.memory_activity_requested = opts_.enable_memory_tracking; input.memory_activity_configured = memory_activity_configured_.load(std::memory_order_acquire); + input.synchronization_requested = opts_.enable_synchronization; + input.synchronization_configured = + synchronization_configured_.load(std::memory_order_acquire); input.kernel_rows = delta(kernel_rows, capability_kernel_rows_baseline_); input.memcpy_rows = delta(memcpy_rows, capability_memcpy_rows_baseline_); + input.synchronization_rows = delta( + synchronization_rows, capability_synchronization_rows_baseline_); input.memory_activity_rows = delta( memory_activity_rows, capability_memory_activity_rows_baseline_); input.profiling_sample_rows = @@ -588,6 +675,7 @@ void RocprofilerBackend::emitCapabilities() { capability_kernel_rows_baseline_ = kernel_rows; capability_memcpy_rows_baseline_ = memcpy_rows; + capability_synchronization_rows_baseline_ = synchronization_rows; capability_memory_activity_rows_baseline_ = memory_activity_rows; capability_pm_sample_rows_baseline_ = pm_sample_rows; capability_dropped_records_baseline_ = dropped_records; @@ -794,6 +882,13 @@ void RocprofilerBackend::bufferTracingShim(rocprofiler_context_id_t, backend->handleMemoryCopy(*record); break; } + case ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API_EXT: { + const auto* record = static_cast< + const rocprofiler_buffer_tracing_hip_api_ext_record_t*>( + header->payload); + backend->handleSynchronization(*record); + break; + } case ROCPROFILER_BUFFER_TRACING_MEMORY_ALLOCATION: { const auto* record = static_cast< const rocprofiler_buffer_tracing_memory_allocation_record_t*>( @@ -989,7 +1084,8 @@ void RocprofilerBackend::handleKernelDispatch( out.type = TraceType::KERNEL; out.device_id = *device_id; out.stream = static_cast(info.queue_id.handle); - out.cpu_start_ns = static_cast(start_timestamp); + out.cpu_start_ns = trace_epoch_offset_ns_ + + static_cast(start_timestamp); out.duration_ns = static_cast(end_timestamp >= start_timestamp ? end_timestamp - start_timestamp : 0); @@ -1139,7 +1235,8 @@ void RocprofilerBackend::handleMemoryCopy( ActivityRecord out{}; out.type = TraceType::MEMCPY; out.device_id = device_id.value_or(0); - out.cpu_start_ns = static_cast(data.start_timestamp); + out.cpu_start_ns = trace_epoch_offset_ns_ + + static_cast(data.start_timestamp); out.duration_ns = static_cast( data.end_timestamp >= data.start_timestamp ? data.end_timestamp - data.start_timestamp : 0); @@ -1171,6 +1268,99 @@ void RocprofilerBackend::handleMemoryCopy( } } +uint32_t RocprofilerBackend::internHipStreamHandle(const uint64_t handle) { + if (handle == 0) return 0; + std::lock_guard lock(synchronization_handle_mutex_); + const auto [itr, inserted] = hip_stream_ids_.emplace(handle, 0); + if (inserted) itr->second = next_hip_stream_id_++; + return itr->second; +} + +uint32_t RocprofilerBackend::internHipEventHandle(const uint64_t handle) { + if (handle == 0) return 0; + std::lock_guard lock(synchronization_handle_mutex_); + const auto [itr, inserted] = hip_event_ids_.emplace(handle, 0); + if (inserted) itr->second = next_hip_event_id_++; + return itr->second; +} + +void RocprofilerBackend::handleSynchronization( + const rocprofiler_buffer_tracing_hip_api_ext_record_t& data) { + const auto operation = ClassifySynchronizationOperation(data.operation); + if (!operation.has_value()) return; + + uint64_t stream_handle = 0; + uint64_t event_handle = 0; + switch (data.operation) { + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipEventSynchronize): + event_handle = reinterpret_cast( + data.args.hipEventSynchronize.event); + break; + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent): + stream_handle = reinterpret_cast( + data.args.hipStreamWaitEvent.stream); + event_handle = reinterpret_cast( + data.args.hipStreamWaitEvent.event); + break; + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent_spt): + stream_handle = reinterpret_cast( + data.args.hipStreamWaitEvent_spt.stream); + event_handle = reinterpret_cast( + data.args.hipStreamWaitEvent_spt.event); + break; + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamSynchronize): + stream_handle = reinterpret_cast( + data.args.hipStreamSynchronize.stream); + break; + case static_cast( + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamSynchronize_spt): + stream_handle = reinterpret_cast( + data.args.hipStreamSynchronize_spt.stream); + break; + default: + break; + } + + ActivityRecord out{}; + out.type = TraceType::SYNCHRONIZATION; + out.cpu_start_ns = trace_epoch_offset_ns_ + + static_cast(data.start_timestamp); + out.duration_ns = static_cast( + data.end_timestamp >= data.start_timestamp + ? data.end_timestamp - data.start_timestamp + : 0); + out.api_start_ns = out.cpu_start_ns; + out.api_exit_ns = out.cpu_start_ns + out.duration_ns; + out.sync_type = ResolveAmdSynchronizationType(*operation); + out.stream = internHipStreamHandle(stream_handle); + out.sync_event_id = internHipEventHandle(event_handle); + out.context_id = 0; + out.corr_id = TruncateCorrelationId(data.correlation_id.internal); + + if (data.correlation_id.external.value != 0) { + std::lock_guard lock(external_scope_mutex_); + if (auto itr = external_scope_metadata_.find( + data.correlation_id.external.value); + itr != external_scope_metadata_.end() && + !itr->second.user_scope.empty()) { + std::snprintf(out.user_scope, sizeof(out.user_scope), "%s", + itr->second.user_scope.c_str()); + out.scope_depth = itr->second.scope_depth; + } + } + + if (g_monitorBuffer.Push(out)) { + synchronization_rows_emitted_.fetch_add(1, + std::memory_order_relaxed); + } else { + trace_records_queue_dropped_.fetch_add(1, std::memory_order_relaxed); + } +} + void RocprofilerBackend::handleMemoryAllocation( const rocprofiler_buffer_tracing_memory_allocation_record_t& data) { const auto operation = NormalizeAmdMemoryAllocationOperation( @@ -1222,7 +1412,8 @@ void RocprofilerBackend::handleMemoryAllocation( ActivityRecord out{}; out.type = TraceType::MEMORY_ALLOC; out.device_id = metadata.device_id; - out.cpu_start_ns = static_cast(data.start_timestamp); + out.cpu_start_ns = trace_epoch_offset_ns_ + + static_cast(data.start_timestamp); out.duration_ns = static_cast( data.end_timestamp >= data.start_timestamp ? data.end_timestamp - data.start_timestamp diff --git a/include/gpufl/backends/amd/rocprofiler_backend.hpp b/include/gpufl/backends/amd/rocprofiler_backend.hpp index 0c5fb72..5ba7765 100644 --- a/include/gpufl/backends/amd/rocprofiler_backend.hpp +++ b/include/gpufl/backends/amd/rocprofiler_backend.hpp @@ -92,11 +92,15 @@ class RocprofilerBackend final : public IMonitorBackend { uint64_t end_timestamp, const rocprofiler_async_correlation_id_t& correlation_id); void handleMemoryCopy(const rocprofiler_buffer_tracing_memory_copy_record_t& data); + void handleSynchronization( + const rocprofiler_buffer_tracing_hip_api_ext_record_t& data); void handleMemoryAllocation( const rocprofiler_buffer_tracing_memory_allocation_record_t& data); void handleCodeObjectLoad(const rocprofiler_callback_tracing_code_object_load_data_t& data); std::string resolveKernelName(uint64_t kernel_id) const; + uint32_t internHipStreamHandle(uint64_t handle); + uint32_t internHipEventHandle(uint64_t handle); AmdTraceEndpoint resolveTraceEndpoint( rocprofiler_agent_id_t agent_id) const; uint32_t classifyMemcpyKind(rocprofiler_agent_id_t src_agent, @@ -131,6 +135,12 @@ class RocprofilerBackend final : public IMonitorBackend { mutable std::mutex memory_allocation_mutex_; std::unordered_map memory_allocations_; + mutable std::mutex synchronization_handle_mutex_; + std::unordered_map hip_stream_ids_; + std::unordered_map hip_event_ids_; + uint32_t next_hip_stream_id_ = 1; + uint32_t next_hip_event_id_ = 1; + mutable std::mutex agent_mutex_; std::unordered_map gpu_device_ids_; std::unordered_map agent_types_; @@ -164,6 +174,7 @@ class RocprofilerBackend final : public IMonitorBackend { std::atomic kernel_rows_emitted_{0}; std::atomic memcpy_rows_emitted_{0}; + std::atomic synchronization_rows_emitted_{0}; std::atomic memory_activity_rows_emitted_{0}; std::atomic trace_records_dropped_{0}; std::atomic trace_records_queue_dropped_{0}; @@ -175,6 +186,7 @@ class RocprofilerBackend final : public IMonitorBackend { mutable std::string capture_capabilities_session_id_; mutable uint64_t capability_kernel_rows_baseline_ = 0; mutable uint64_t capability_memcpy_rows_baseline_ = 0; + mutable uint64_t capability_synchronization_rows_baseline_ = 0; mutable uint64_t capability_memory_activity_rows_baseline_ = 0; mutable uint64_t capability_pm_sample_rows_baseline_ = 0; mutable uint64_t capability_dropped_records_baseline_ = 0; @@ -183,6 +195,8 @@ class RocprofilerBackend final : public IMonitorBackend { mutable uint64_t capability_unattributed_records_baseline_ = 0; mutable uint64_t capability_scope_correlation_failures_baseline_ = 0; + // Fixed before tracing starts; converts ROCprofiler's clock to Unix ns. + int64_t trace_epoch_offset_ns_ = 0; std::atomic initialized_{false}; std::atomic active_{false}; std::atomic start_requested_{false}; @@ -191,6 +205,7 @@ class RocprofilerBackend final : public IMonitorBackend { std::atomic start_failure_logged_{false}; std::mutex start_stop_mutex_; std::atomic tool_registered_{false}; + std::atomic synchronization_configured_{false}; std::atomic memory_activity_configured_{false}; }; diff --git a/include/gpufl/core/dictionary_manager.cpp b/include/gpufl/core/dictionary_manager.cpp index 4cedbee..440e8d6 100644 --- a/include/gpufl/core/dictionary_manager.cpp +++ b/include/gpufl/core/dictionary_manager.cpp @@ -874,7 +874,8 @@ void DictionaryManager::flushDisassembly(Logger& logger, } poss << "]}"; if (!pfirst) { - logger.write(DictLine{poss.str()}); + // Static mappings are data, not dictionaries: emit only once. + logger.write(SassLine{poss.str()}); } } } diff --git a/include/gpufl/core/monitor.cpp b/include/gpufl/core/monitor.cpp index 86c0f5a..eb6ec35 100644 --- a/include/gpufl/core/monitor.cpp +++ b/include/gpufl/core/monitor.cpp @@ -880,6 +880,10 @@ uint64_t Monitor::PmSampleRowsSeen() { return g_state.batches.pmSampleRowsSeen(); } +uint64_t Monitor::SynchronizationRowsSeen() { + return g_state.batches.synchronizationRowsSeen(); +} + uint64_t Monitor::MemoryAllocRowsSeen() { return g_state.batches.memoryAllocRowsSeen(); } diff --git a/include/gpufl/core/monitor.hpp b/include/gpufl/core/monitor.hpp index 7e9cc4d..9975371 100644 --- a/include/gpufl/core/monitor.hpp +++ b/include/gpufl/core/monitor.hpp @@ -172,10 +172,10 @@ struct MonitorOptions { // InitOptions::enable_external_correlation; copied across in the // gpufl::init() → CuptiBackend::initialize() conversion path. bool enable_external_correlation = true; - // Gate for CUPTI_ACTIVITY_KIND_SYNCHRONIZATION. Mirror of - // InitOptions::enable_synchronization. Backend honors this in - // CuptiBackend::start() - flag false means we never call - // cuptiActivityEnable for the kind, so zero overhead. + // Gate for synchronization activity. NVIDIA uses CUPTI activity; + // AMD trace sessions use filtered ROCprofiler HIP runtime API records. + // Mirror of InitOptions::enable_synchronization. False leaves the + // vendor service disabled, so there is no synchronization record volume. bool enable_synchronization = true; // Gate for CUPTI_ACTIVITY_KIND_MEMORY2. bool enable_memory_tracking = true; @@ -458,6 +458,8 @@ class Monitor { static uint64_t ScopeAttributionTruncated(); /** @brief PM metric rows that passed through scope attribution. */ static uint64_t PmSampleRowsSeen(); + /** @brief Synchronization rows accepted by the batch pipeline. */ + static uint64_t SynchronizationRowsSeen(); /** @brief Memory-allocation rows accepted by the batch pipeline. */ static uint64_t MemoryAllocRowsSeen(); diff --git a/include/gpufl/core/monitor_batch_manager.cpp b/include/gpufl/core/monitor_batch_manager.cpp index fafb3c2..54688b6 100644 --- a/include/gpufl/core/monitor_batch_manager.cpp +++ b/include/gpufl/core/monitor_batch_manager.cpp @@ -41,6 +41,7 @@ void MonitorBatchManager::reset() { } syncBatch_.clear(); memAllocBatch_.clear(); + synchronizationRowsSeen_.store(0, std::memory_order_relaxed); memoryAllocRowsSeen_.store(0, std::memory_order_relaxed); pendingDetails_.clear(); @@ -433,6 +434,10 @@ uint64_t MonitorBatchManager::pmSampleRowsSeen() const { return pmSampleRowsSeen_; } +uint64_t MonitorBatchManager::synchronizationRowsSeen() const { + return synchronizationRowsSeen_.load(std::memory_order_relaxed); +} + uint64_t MonitorBatchManager::memoryAllocRowsSeen() const { return memoryAllocRowsSeen_.load(std::memory_order_relaxed); } @@ -615,6 +620,7 @@ bool MonitorBatchManager::pushMemoryAlloc(const MemoryAllocEventBatchRow& row) { } void MonitorBatchManager::pushSynchronization(const SynchronizationEventBatchRow& row) { + synchronizationRowsSeen_.fetch_add(1, std::memory_order_relaxed); syncBatch_.push(row); } diff --git a/include/gpufl/core/monitor_batch_manager.hpp b/include/gpufl/core/monitor_batch_manager.hpp index 4d0e4f3..fe950d0 100644 --- a/include/gpufl/core/monitor_batch_manager.hpp +++ b/include/gpufl/core/monitor_batch_manager.hpp @@ -84,6 +84,8 @@ class MonitorBatchManager { uint64_t scopeAttributionTruncated() const; /** @brief PM metric rows that have passed through scope attribution. */ uint64_t pmSampleRowsSeen() const; + /** @brief Synchronization rows accepted by the batch pipeline. */ + uint64_t synchronizationRowsSeen() const; /** @brief Memory-allocation rows accepted by the batch pipeline. */ uint64_t memoryAllocRowsSeen() const; @@ -227,6 +229,7 @@ class MonitorBatchManager { uint64_t pmSampleRowsSeen_ = 0; BatchBuffer syncBatch_; + std::atomic synchronizationRowsSeen_{0}; std::atomic memoryAllocRowsSeen_{0}; BatchBuffer memAllocBatch_; uint64_t syncBatchId_ = 0; diff --git a/include/gpufl/gpufl.hpp b/include/gpufl/gpufl.hpp index ac37f0e..573b7cd 100644 --- a/include/gpufl/gpufl.hpp +++ b/include/gpufl/gpufl.hpp @@ -63,16 +63,16 @@ struct InitOptions { // disable only if running on a CUPTI version that errors on the // kind (logs a soft-warning if so, doesn't crash). bool enable_external_correlation = true; - // Enable CUPTI_ACTIVITY_KIND_SYNCHRONIZATION so that every - // cudaStreamSynchronize / cudaDeviceSynchronize / cudaEventSynchronize - // / cuStreamWaitEvent call gets a wall-clock timed record. The - // primary insight: time spent here = host blocked on GPU = a - // direct measure of GPU underutilization. Default-on because the - // overhead is small (one record per sync call, mid volume) and - // the answer it unlocks ("X% of your wall time is `cudaStreamSync`") - // is a top-five most-asked question. If a workload performs - // millions of synchronizations and the volume becomes a problem, - // disable this flag - the rest of the pipeline keeps working. + // Enable synchronization activity. NVIDIA uses CUPTI synchronization + // records; AMD trace sessions use filtered ROCprofiler HIP runtime API + // records. Event synchronize, stream wait-event, stream synchronize, + // and context/device synchronize calls are normalized to portable wire + // values with host-observed API durations. Most synchronize calls expose + // host time blocked on GPU work; stream wait-event may only enqueue a + // dependency. Default-on because the overhead is small (one record per + // selected call, mid volume). If a workload performs millions of + // synchronizations and the volume becomes a problem, disable this flag; + // the rest of the pipeline keeps working. bool enable_synchronization = true; // Enable CUPTI_ACTIVITY_KIND_MEMORY2 to capture cudaMalloc / // cudaFree / cudaMallocAsync / cudaMallocManaged / cudaMallocHost diff --git a/include/gpufl/report/text_report.cpp b/include/gpufl/report/text_report.cpp index 565e46d..c7e768b 100644 --- a/include/gpufl/report/text_report.cpp +++ b/include/gpufl/report/text_report.cpp @@ -178,6 +178,16 @@ std::string resolveMemoryKind(int kind) { } } +std::string resolveSynchronizationType(int type) { + switch (type) { + case 1: return "Event Synchronize"; + case 2: return "Stream Wait Event"; + case 3: return "Stream Synchronize"; + case 4: return "Context Synchronize"; + default: return "Unknown(" + std::to_string(type) + ")"; + } +} + // Values match CUpti_ActivityPCSamplingStallReason enum from cupti_activity.h const std::map kStallNames = { {2, "Instruction Fetch"}, {3, "Execution Dependency"}, @@ -411,6 +421,39 @@ bool TextReport::isAmdSession() const { identity.find("advanced micro devices") != std::string::npos; } +bool TextReport::tryParseSynchronizationRecord(const JsonValue& rec) { + const std::string type = rec.value("type", ""); + if (type == "synchronization_event_batch") { + auto ci = buildColumnIndex(rec["columns"]); + const int64_t base = rec.value("base_time_ns", 0); + for (const auto& row : rec["rows"].get_array()) { + const int64_t duration_ns = rowInt(row, ci, "duration_ns"); + synchronizations_.push_back({ + base + rowInt(row, ci, "dt_ns"), + duration_ns / 1e6, + static_cast(rowInt(row, ci, "sync_type")), + static_cast(rowInt(row, ci, "stream_id")), + static_cast(rowInt(row, ci, "event_id")), + static_cast(rowInt(row, ci, "context_id")), + }); + } + return true; + } + + if (type == "synchronization_event") { + synchronizations_.push_back({ + rec.value("start_ns", 0), + rec.value("duration_ns", 0) / 1e6, + static_cast(rec.value("sync_type", 0)), + rec.value("stream_id", 0), + rec.value("event_id", 0), + rec.value("context_id", 0), + }); + return true; + } + return false; +} + bool TextReport::tryParseMemoryAllocationRecord(const JsonValue& rec) { const std::string type = rec.value("type", ""); if (type == "memory_alloc_event_batch") { @@ -447,6 +490,7 @@ void TextReport::parseDeviceLog(const std::vector& records, for (const auto& rec : records) { const std::string type = rec.value("type", ""); + if (tryParseSynchronizationRecord(rec)) continue; if (tryParseMemoryAllocationRecord(rec)) continue; if ((type == "job_start" || type == "init") && info_.app_name.empty()) { @@ -551,6 +595,7 @@ void TextReport::parseScopeLog(const std::vector& records, for (const auto& rec : records) { const std::string type = rec.value("type", ""); + if (tryParseSynchronizationRecord(rec)) continue; if (tryParseMemoryAllocationRecord(rec)) continue; if ((type == "job_start" || type == "init") && info_.app_name.empty()) { @@ -688,6 +733,7 @@ std::string TextReport::generate() const { writeKernelDetails(out); writeMemcpySummary(out); writeMemoryAllocationSummary(out); + writeSynchronizationSummary(out); writeSystemMetrics(out); writeScopeSummary(out); writePerfMetricsSummary(out); @@ -1053,6 +1099,45 @@ void TextReport::writeMemoryAllocationSummary(std::ostringstream& out) const { << " Runtime and allocator-internal activity may be included.\n"; } +void TextReport::writeSynchronizationSummary(std::ostringstream& out) const { + out << "\n" << SEP << "\n Synchronization Summary\n" << SEP << "\n"; + if (synchronizations_.empty()) { + out << " (No synchronization data)\n"; + return; + } + + AggStats total; + std::map grouped; + for (const auto& record : synchronizations_) { + total.add(record.duration_ms); + grouped[record.sync_type].add(record.duration_ms); + } + + out << " Total Calls: " << synchronizations_.size() << "\n"; + out << " Total API Time: " << fmtDuration(total.total) << "\n"; + out << " Avg API Time: " << fmtDuration(total.avg()) << "\n"; + out << " Max API Time: " << fmtDuration(total.max_val) << "\n\n"; + + out << " By Synchronization Type:\n"; + out << " " << std::left << std::setw(28) << "Type" + << std::right << std::setw(9) << "Calls" + << std::setw(14) << "Total" + << std::setw(14) << "Avg" + << std::setw(14) << "Max" << "\n"; + out << " " << std::string(65, '-') << "\n"; + for (const auto& [type, stats] : grouped) { + out << " " << std::left << std::setw(28) + << resolveSynchronizationType(type) + << std::right << std::setw(9) << stats.count + << std::setw(14) << fmtDuration(stats.total) + << std::setw(14) << fmtDuration(stats.avg()) + << std::setw(14) << fmtDuration(stats.max_val) << "\n"; + } + + out << "\n Note: API time is the host-observed HIP/CUDA call duration.\n" + << " Stream Wait Event may enqueue a dependency without blocking.\n"; +} + void TextReport::writeSystemMetrics(std::ostringstream& out) const { out << "\n" << SEP << "\n System Metrics\n" << SEP << "\n"; if (device_metrics_.empty() && host_metrics_.empty()) { diff --git a/include/gpufl/report/text_report.hpp b/include/gpufl/report/text_report.hpp index ee3fca0..16018ce 100644 --- a/include/gpufl/report/text_report.hpp +++ b/include/gpufl/report/text_report.hpp @@ -71,6 +71,15 @@ class TextReport { uint64_t bytes = 0; }; + struct SynchronizationRecord { + int64_t start_ns = 0; + double duration_ms = 0; + uint8_t sync_type = 0; + uint32_t stream_id = 0; + uint32_t event_id = 0; + uint32_t context_id = 0; + }; + struct DeviceMetricRecord { int64_t ts_ns = 0; int gpu_util = 0; @@ -180,6 +189,7 @@ class TextReport { std::vector kernels_; std::vector memcpy_; std::vector memory_allocations_; + std::vector synchronizations_; std::vector device_metrics_; std::vector host_metrics_; std::vector scope_events_; @@ -202,6 +212,7 @@ class TextReport { std::unordered_map& metric_dict); void parseDeviceLog(const std::vector& records, const std::unordered_map& kernel_dict); + bool tryParseSynchronizationRecord(const JsonValue& record); bool tryParseMemoryAllocationRecord(const JsonValue& record); void parseScopeLog(const std::vector& records, const std::unordered_map& scope_name_dict, @@ -223,6 +234,7 @@ class TextReport { void writeKernelDetails(std::ostringstream& out) const; void writeMemcpySummary(std::ostringstream& out) const; void writeMemoryAllocationSummary(std::ostringstream& out) const; + void writeSynchronizationSummary(std::ostringstream& out) const; void writeSystemMetrics(std::ostringstream& out) const; void writeScopeSummary(std::ostringstream& out) const; void writePerfMetricsSummary(std::ostringstream& out) const; diff --git a/tests/backends/amd/test_amd_profiling_policy.cpp b/tests/backends/amd/test_amd_profiling_policy.cpp index 36023c0..4530eda 100644 --- a/tests/backends/amd/test_amd_profiling_policy.cpp +++ b/tests/backends/amd/test_amd_profiling_policy.cpp @@ -340,3 +340,41 @@ TEST(AmdCaptureCapabilities, MemoryActivityReportsCollectionAndAvailability) { ASSERT_NE(memory, nullptr); EXPECT_EQ(memory->status, "not_requested"); } + +TEST(AmdCaptureCapabilities, SynchronizationReportsCollectionAndAvailability) { + gpufl::amd::AmdCaptureCapabilityInput input; + input.session_id = "amd-session"; + input.plan = gpufl::amd::ResolveAmdProfilingPlan( + gpufl::ProfilingEngine::Trace, {}); + input.trace_configured = true; + input.synchronization_requested = true; + input.synchronization_configured = true; + input.synchronization_rows = 4; + + auto event = gpufl::amd::BuildAmdCaptureCapabilitiesEvent(input); + const auto* sync = FindCapability(event, "sync_activity"); + ASSERT_NE(sync, nullptr); + EXPECT_TRUE(sync->requested); + EXPECT_EQ(sync->status, "collected"); + EXPECT_EQ(sync->mode, "rocprofiler_hip_runtime_api_ext"); + + input.synchronization_rows = 0; + event = gpufl::amd::BuildAmdCaptureCapabilitiesEvent(input); + sync = FindCapability(event, "sync_activity"); + ASSERT_NE(sync, nullptr); + EXPECT_EQ(sync->status, "enabled_no_data"); + EXPECT_EQ(sync->reason_code, "enabled_but_no_records"); + + input.synchronization_configured = false; + event = gpufl::amd::BuildAmdCaptureCapabilitiesEvent(input); + sync = FindCapability(event, "sync_activity"); + ASSERT_NE(sync, nullptr); + EXPECT_EQ(sync->status, "skipped"); + EXPECT_EQ(sync->reason_code, "rocprofiler_hip_runtime_api_unavailable"); + + input.synchronization_requested = false; + event = gpufl::amd::BuildAmdCaptureCapabilitiesEvent(input); + sync = FindCapability(event, "sync_activity"); + ASSERT_NE(sync, nullptr); + EXPECT_EQ(sync->status, "not_requested"); +} diff --git a/tests/backends/amd/test_amd_trace_policy.cpp b/tests/backends/amd/test_amd_trace_policy.cpp index 00dabba..5f46574 100644 --- a/tests/backends/amd/test_amd_trace_policy.cpp +++ b/tests/backends/amd/test_amd_trace_policy.cpp @@ -67,3 +67,19 @@ TEST(AmdTracePolicy, MemoryAllocationKindIsTruthfulAcrossAgents) { EXPECT_EQ(gpufl::amd::ResolveAmdMemoryAllocationKind(Cpu()), 0u); EXPECT_EQ(gpufl::amd::ResolveAmdMemoryAllocationKind({}), 0u); } + +TEST(AmdTracePolicy, SynchronizationOperationsUsePortableWireValues) { + using Operation = gpufl::amd::AmdSynchronizationOperation; + EXPECT_EQ(gpufl::amd::ResolveAmdSynchronizationType( + Operation::EventSynchronize), + 1u); + EXPECT_EQ(gpufl::amd::ResolveAmdSynchronizationType( + Operation::StreamWaitEvent), + 2u); + EXPECT_EQ(gpufl::amd::ResolveAmdSynchronizationType( + Operation::StreamSynchronize), + 3u); + EXPECT_EQ(gpufl::amd::ResolveAmdSynchronizationType( + Operation::ContextSynchronize), + 4u); +} diff --git a/tests/core/test_monitor.cpp b/tests/core/test_monitor.cpp index 5ed22f8..8a597dd 100644 --- a/tests/core/test_monitor.cpp +++ b/tests/core/test_monitor.cpp @@ -523,3 +523,21 @@ TEST(MemoryAllocationBatchTest, CountsAcceptedRowsAndResets) { manager.reset(); EXPECT_EQ(manager.memoryAllocRowsSeen(), 0u); } + +TEST(SynchronizationBatchTest, CountsAcceptedRowsAndResets) { + gpufl::detail::MonitorBatchManager manager; + manager.reset(); + + EXPECT_EQ(manager.synchronizationRowsSeen(), 0u); + gpufl::SynchronizationEventBatchRow row{}; + row.start_ns = 100; + manager.pushSynchronization(row); + EXPECT_EQ(manager.synchronizationRowsSeen(), 1u); + + row.start_ns = 200; + manager.pushSynchronization(row); + EXPECT_EQ(manager.synchronizationRowsSeen(), 2u); + + manager.reset(); + EXPECT_EQ(manager.synchronizationRowsSeen(), 0u); +} diff --git a/tests/core/test_text_report.cpp b/tests/core/test_text_report.cpp index 97574c4..50edc6d 100644 --- a/tests/core/test_text_report.cpp +++ b/tests/core/test_text_report.cpp @@ -132,4 +132,42 @@ TEST_F(TextReportTest, LegacyMemoryAllocationEventProducesSummary) { EXPECT_NE(report.find("Managed"), std::string::npos); } +TEST_F(TextReportTest, SynchronizationBatchRowsProduceSummary) { + WriteLog("scope", { + R"({"type":"job_start","session_id":"s1","app":"sync_report","ts_ns":1000,"gpu_static_devices":[{"name":"AMD Radeon Test","vendor":"AMD","multi_processor_count":32}]})", + R"({"type":"synchronization_event_batch","session_id":"s1","base_time_ns":1000,"columns":["dt_ns","duration_ns","sync_type","stream_id","event_id","context_id","corr_id","function_id"],"rows":[[10,1000,1,0,1,0,1,0],[20,2000,2,1,1,0,2,0],[30,3000,3,1,0,0,3,0],[40,4000,4,0,0,0,4,0]]})", + R"({"type":"shutdown","session_id":"s1","ts_ns":10000})", + }); + + const std::string report = Generate(); + EXPECT_NE(report.find("Synchronization Summary"), std::string::npos); + EXPECT_NE(report.find("Total Calls: 4"), std::string::npos); + EXPECT_NE(report.find("Total API Time: 10.00 usec"), std::string::npos); + EXPECT_NE(report.find("Avg API Time: 2.50 usec"), std::string::npos); + EXPECT_NE(report.find("Max API Time: 4.00 usec"), std::string::npos); + EXPECT_NE(report.find("Event Synchronize"), std::string::npos); + EXPECT_NE(report.find("Stream Wait Event"), std::string::npos); + EXPECT_NE(report.find("Stream Synchronize"), std::string::npos); + EXPECT_NE(report.find("Context Synchronize"), std::string::npos); + + const auto allocations = report.find("Memory Allocation Summary"); + const auto synchronization = report.find("Synchronization Summary"); + const auto system_metrics = report.find("System Metrics"); + EXPECT_LT(allocations, synchronization); + EXPECT_LT(synchronization, system_metrics); +} + +TEST_F(TextReportTest, LegacySynchronizationEventProducesSummary) { + WriteLog("device", { + R"({"type":"job_start","session_id":"s1","app":"legacy_sync_report","ts_ns":1000,"gpu_static_devices":[{"name":"NVIDIA Test GPU","vendor":"NVIDIA","multi_processor_count":10}]})", + R"({"type":"synchronization_event","session_id":"s1","start_ns":1100,"end_ns":6100,"duration_ns":5000,"sync_type":1,"stream_id":0,"event_id":7,"context_id":0,"corr_id":1})", + R"({"type":"shutdown","session_id":"s1","ts_ns":10000})", + }); + + const std::string report = Generate(); + EXPECT_NE(report.find("Total Calls: 1"), std::string::npos); + EXPECT_NE(report.find("Total API Time: 5.00 usec"), std::string::npos); + EXPECT_NE(report.find("Event Synchronize"), std::string::npos); +} + } // namespace