From cf696f04be7587a68d95e724fbe006363cafe0bf Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Tue, 1 Sep 2026 22:55:34 -0700 Subject: [PATCH] feat(amd): add memory allocation tracing --- example/amd/CMakeLists.txt | 4 + example/amd/README.md | 19 ++- example/amd/memory_allocation_rows.cpp | 152 ++++++++++++++++++ .../backends/amd/amd_capture_capabilities.cpp | 24 +++ .../backends/amd/amd_capture_capabilities.hpp | 3 + .../gpufl/backends/amd/amd_trace_policy.cpp | 19 +++ .../gpufl/backends/amd/amd_trace_policy.hpp | 17 ++ .../backends/amd/rocprofiler_backend.cpp | 120 ++++++++++++++ .../backends/amd/rocprofiler_backend.hpp | 13 ++ include/gpufl/core/activity_record.hpp | 18 +-- include/gpufl/core/events/memory_events.hpp | 43 ++--- include/gpufl/core/model/batch_models.cpp | 4 +- include/gpufl/core/monitor.cpp | 4 + include/gpufl/core/monitor.hpp | 2 + include/gpufl/core/monitor_batch_manager.cpp | 6 + include/gpufl/core/monitor_batch_manager.hpp | 3 + include/gpufl/core/trace_type.hpp | 26 ++- include/gpufl/report/text_report.cpp | 137 ++++++++++++++++ include/gpufl/report/text_report.hpp | 11 ++ .../amd/test_amd_profiling_policy.cpp | 39 +++++ tests/backends/amd/test_amd_trace_policy.cpp | 28 ++++ tests/core/test_monitor.cpp | 18 +++ tests/core/test_text_report.cpp | 39 +++++ 23 files changed, 693 insertions(+), 56 deletions(-) create mode 100644 example/amd/memory_allocation_rows.cpp diff --git a/example/amd/CMakeLists.txt b/example/amd/CMakeLists.txt index f8d0d49..91b5dfe 100644 --- a/example/amd/CMakeLists.txt +++ b/example/amd/CMakeLists.txt @@ -48,6 +48,7 @@ list(APPEND HIP_HIPCC_FLAGS "-g") set(AMD_EXAMPLE_SOURCES check_device.cpp gpufl_scope_demo.cpp + memory_allocation_rows.cpp pm_sampling_sample_rows.cpp vector_add_benchmark.cpp ) @@ -59,17 +60,20 @@ set_source_files_properties( 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) 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) 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) diff --git a/example/amd/README.md b/example/amd/README.md index 4a5db05..0b0878b 100644 --- a/example/amd/README.md +++ b/example/amd/README.md @@ -8,6 +8,7 @@ This folder mirrors the CUDA example area with runnable HIP examples for AMD GPU - AMD static device inventory via HIP - AMD kernel dispatch tracing via `rocprofiler-sdk` - AMD memcpy tracing via `rocprofiler-sdk` +- AMD memory-allocation tracing via `rocprofiler-sdk` - Per-dispatch AMD hardware counters via ROCprofiler dispatch counting - Device-wide `PmSampling` timelines via ROCprofiler device counting - `gpufl` initialization with `backend = gpufl::BackendKind::Amd` @@ -23,7 +24,7 @@ Today, the AMD backend is useful for: - system metric logging - device inventory -- automatic HIP kernel and memcpy tracing +- automatic HIP kernel, memcpy, and memory-allocation tracing - per-dispatch and device-wide hardware-counter profiling - scope-level application instrumentation @@ -39,6 +40,8 @@ It is not yet useful for: - HIP vector add benchmark with 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 @@ -57,6 +60,7 @@ cmake -S . -B build-rocm-examples \ 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_gpufl_scope_demo +cmake --build build-rocm-examples --target amd_memory_allocation_rows cmake --build build-rocm-examples --target amd_pm_sampling_sample_rows ``` @@ -113,9 +117,15 @@ subproject and disables the parent example/test targets to avoid recursion. ./build-rocm-examples/example/amd/amd_check_device ./build-rocm-examples/example/amd/amd_vector_add_benchmark ./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 ``` +`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_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 @@ -149,6 +159,12 @@ Success! Device 0: AMD Radeon RX 9070 XT (arch gfx1201, capability 12.0) gfl_amd_scope ``` +`amd_memory_allocation_rows` writes logs with prefix: + +```bash +gfl_amd_memory_rows +``` + `amd_pm_sampling_sample_rows` writes logs with prefix: ```bash @@ -162,6 +178,7 @@ With `rocprofiler-sdk` available, expect: - `kernel_event_batch` - `kernel_detail` - `memcpy_event_batch` +- `memory_alloc_event_batch` when `enable_memory_tracking` 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/memory_allocation_rows.cpp b/example/amd/memory_allocation_rows.cpp new file mode 100644 index 0000000..8576309 --- /dev/null +++ b/example/amd/memory_allocation_rows.cpp @@ -0,0 +1,152 @@ +#include + +#include +#include +#include +#include +#include + +#include "gpufl/core/monitor.hpp" +#include "gpufl/gpufl.hpp" + +namespace { + +constexpr int kAllocationsPerPhase = 4; +constexpr uint64_t kExpectedRowsPerPhase = 2 * kAllocationsPerPhase; +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; +} + +bool RunAllocationPhase(const size_t base_bytes) { + std::vector allocations; + allocations.reserve(kAllocationsPerPhase); + + bool ok = true; + for (int index = 0; index < kAllocationsPerPhase; ++index) { + void* allocation = nullptr; + const size_t bytes = base_bytes * static_cast(index + 1); + if (!CheckHip(hipMalloc(&allocation, bytes), "hipMalloc")) { + ok = false; + break; + } + allocations.push_back(allocation); + } + + for (auto it = allocations.rbegin(); it != allocations.rend(); ++it) { + if (!CheckHip(hipFree(*it), "hipFree")) ok = false; + } + return ok; +} + +uint64_t WaitForAllocationRows(const uint64_t minimum_rows) { + const auto deadline = std::chrono::steady_clock::now() + kDeliveryTimeout; + uint64_t rows = gpufl::Monitor::MemoryAllocRowsSeen(); + while (rows < minimum_rows && std::chrono::steady_clock::now() < deadline) { + std::this_thread::sleep_for(std::chrono::milliseconds(10)); + rows = gpufl::Monitor::MemoryAllocRowsSeen(); + } + return rows; +} + +} // namespace + +int main() { + gpufl::InitOptions opts; + opts.app_name = "amd_memory_allocation_rows"; + opts.log_path = "gfl_amd_memory_rows"; + opts.backend = gpufl::BackendKind::Amd; + opts.profiling_engine = gpufl::ProfilingEngine::Trace; + opts.enable_memory_tracking = true; + 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 allocation tracing\n"; + return 1; + } + + const std::string engine = + gpufl::Monitor::ResolvedProfilingEngineWireName(); + const bool engine_ok = engine == "amd.buffer_tracing"; + std::cout << "=== GPUFL AMD Memory Allocation Rows ===\n" + << "Resolved engine: " << engine << "\n"; + if (!engine_ok) { + std::cerr << "Expected amd.buffer_tracing; ROCprofiler tracing may be unavailable\n"; + } + + // The first HIP call loads HSA. GPUFlight then retries its deferred + // ROCprofiler context start from the collector thread. + void* warmup = nullptr; + bool workload_ok = CheckHip(hipMalloc(&warmup, 4096), "warmup hipMalloc"); + if (warmup != nullptr) { + workload_ok = CheckHip(hipFree(warmup), "warmup hipFree") && workload_ok; + } + std::this_thread::sleep_for(std::chrono::milliseconds(50)); + + bool priming_ok = false; + uint64_t rows_before = gpufl::Monitor::MemoryAllocRowsSeen(); + if (workload_ok) { + priming_ok = RunAllocationPhase(16 * 1024); + if (priming_ok) { + rows_before = WaitForAllocationRows(kExpectedRowsPerPhase); + std::this_thread::sleep_for(std::chrono::milliseconds(20)); + rows_before = gpufl::Monitor::MemoryAllocRowsSeen(); + } + } + const bool priming_rows = rows_before >= kExpectedRowsPerPhase; + if (!priming_rows) { + std::cerr << "Priming phase did not emit allocation and free rows\n"; + } + workload_ok = workload_ok && priming_ok && priming_rows; + + bool phase_a_ok = false; + if (workload_ok) { + GFL_SCOPE("memory_rows_phase_a") { + phase_a_ok = RunAllocationPhase(64 * 1024); + } + } + const uint64_t rows_after_a = + WaitForAllocationRows(rows_before + kExpectedRowsPerPhase); + + bool phase_b_ok = false; + if (workload_ok && phase_a_ok) { + GFL_SCOPE("memory_rows_phase_b") { + phase_b_ok = RunAllocationPhase(128 * 1024); + } + } + const uint64_t rows_after_b = + WaitForAllocationRows(rows_after_a + kExpectedRowsPerPhase); + + std::cout << "Allocation 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 allocation rows\n"; + } + if (!phase_b_rows) { + std::cerr << "Phase B did not emit allocation rows\n"; + } + + 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 AMD memory-allocation rows.\n" + << "Inspect logs with prefix " << opts.log_path + << " for memory_alloc_event_batch events.\n"; + return 0; +} diff --git a/include/gpufl/backends/amd/amd_capture_capabilities.cpp b/include/gpufl/backends/amd/amd_capture_capabilities.cpp index a224a99..16bf857 100644 --- a/include/gpufl/backends/amd/amd_capture_capabilities.cpp +++ b/include/gpufl/backends/amd/amd_capture_capabilities.cpp @@ -89,6 +89,30 @@ CaptureCapabilitiesEvent BuildAmdCaptureCapabilitiesEvent( ? "Memory-copy tracing was enabled but emitted no rows this session." : "Memory-copy tracing was not active.")); + AddCapability( + event, "memory_activity", input.memory_activity_requested, + !input.memory_activity_requested + ? "not_requested" + : (input.memory_activity_configured + ? (input.memory_activity_rows > 0 ? "collected" + : "enabled_no_data") + : "skipped"), + input.memory_activity_configured ? "rocprofiler_memory_allocation" + : "disabled", + input.memory_activity_configured + ? (input.memory_activity_rows > 0 ? "" : "enabled_but_no_records") + : (input.memory_activity_requested + ? "rocprofiler_memory_allocation_unavailable" + : ""), + input.memory_activity_configured + ? (input.memory_activity_rows > 0 + ? "ROCprofiler memory-allocation records were collected." + : "ROCprofiler memory-allocation tracing was enabled but " + "emitted no rows this session.") + : (input.memory_activity_requested + ? "ROCprofiler memory-allocation tracing was requested but unavailable." + : "Memory-allocation tracing was not requested.")); + AddCapability( event, "dispatch_counting", dispatch_requested, !dispatch_requested diff --git a/include/gpufl/backends/amd/amd_capture_capabilities.hpp b/include/gpufl/backends/amd/amd_capture_capabilities.hpp index cce4b01..a9ad61d 100644 --- a/include/gpufl/backends/amd/amd_capture_capabilities.hpp +++ b/include/gpufl/backends/amd/amd_capture_capabilities.hpp @@ -12,8 +12,11 @@ struct AmdCaptureCapabilityInput { int64_t ts_ns = 0; AmdResolvedProfilingPlan plan; bool trace_configured = false; + bool memory_activity_requested = false; + bool memory_activity_configured = false; uint64_t kernel_rows = 0; uint64_t memcpy_rows = 0; + uint64_t memory_activity_rows = 0; uint64_t profiling_sample_rows = 0; uint64_t pm_sample_rows = 0; uint64_t dropped_trace_records = 0; diff --git a/include/gpufl/backends/amd/amd_trace_policy.cpp b/include/gpufl/backends/amd/amd_trace_policy.cpp index 13a12bf..11a8aa3 100644 --- a/include/gpufl/backends/amd/amd_trace_policy.cpp +++ b/include/gpufl/backends/amd/amd_trace_policy.cpp @@ -21,4 +21,23 @@ std::optional ResolveAmdMemoryCopyDeviceId( return std::nullopt; } +std::optional ResolveAmdMemoryAllocationDeviceId( + const AmdTraceEndpoint& agent) { + if (agent.kind != AmdTraceAgentKind::Gpu) return std::nullopt; + return agent.device_id; +} + +std::optional NormalizeAmdMemoryAllocationOperation( + const uint32_t operation) { + // rocprofiler_memory_allocation_operation_t: + // 1=ALLOCATE, 2=VMEM_ALLOCATE, 3=FREE, 4=VMEM_FREE. + if (operation == 1 || operation == 2) return uint8_t{1}; + if (operation == 3 || operation == 4) return uint8_t{2}; + return std::nullopt; +} + +uint8_t ResolveAmdMemoryAllocationKind(const AmdTraceEndpoint& agent) { + return agent.kind == AmdTraceAgentKind::Gpu ? uint8_t{3} : uint8_t{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 d0d5733..bc6a9a8 100644 --- a/include/gpufl/backends/amd/amd_trace_policy.hpp +++ b/include/gpufl/backends/amd/amd_trace_policy.hpp @@ -26,4 +26,21 @@ std::optional ResolveAmdMemoryCopyDeviceId( const AmdTraceEndpoint& source, const AmdTraceEndpoint& destination); +// Memory-allocation activity belongs to the agent that owns the allocation. +// CPU allocations intentionally have no GPU device id; callers may still emit +// them as host allocations. An unmapped GPU remains unattributed. +std::optional ResolveAmdMemoryAllocationDeviceId( + const AmdTraceEndpoint& agent); + +// Normalize ROCprofiler's ALLOCATE/VMEM_ALLOCATE/FREE/VMEM_FREE operation +// values to GPUFlight's portable 1=ALLOC, 2=FREE wire contract. NONE and +// unknown future values are not emitted. +std::optional NormalizeAmdMemoryAllocationOperation( + uint32_t operation); + +// GPUFlight memory-kind wire values follow CUPTI for cross-vendor consumers. +// ROCprofiler only identifies the owning agent, so GPU allocations can be +// classified as DEVICE (3) while CPU allocations remain UNKNOWN (0). +uint8_t ResolveAmdMemoryAllocationKind(const AmdTraceEndpoint& agent); + } // namespace gpufl::amd diff --git a/include/gpufl/backends/amd/rocprofiler_backend.cpp b/include/gpufl/backends/amd/rocprofiler_backend.cpp index 169c6c4..e4563ce 100644 --- a/include/gpufl/backends/amd/rocprofiler_backend.cpp +++ b/include/gpufl/backends/amd/rocprofiler_backend.cpp @@ -121,6 +121,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); + 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); trace_buffer_flush_failures_.store(0, std::memory_order_relaxed); @@ -132,6 +133,7 @@ void RocprofilerBackend::initialize(const MonitorOptions& opts) { capture_capabilities_session_id_.clear(); capability_kernel_rows_baseline_ = 0; capability_memcpy_rows_baseline_ = 0; + capability_memory_activity_rows_baseline_ = 0; capability_pm_sample_rows_baseline_ = 0; capability_dropped_records_baseline_ = 0; capability_queue_dropped_records_baseline_ = 0; @@ -166,6 +168,7 @@ void RocprofilerBackend::resetToolState() { client_handle_ = 0; client_finalize_ = nullptr; tool_registered_.store(false); + memory_activity_configured_.store(false); active_.store(false); start_requested_.store(false); next_start_retry_ns_.store(0); @@ -179,6 +182,10 @@ void RocprofilerBackend::resetToolState() { std::lock_guard lock(external_scope_mutex_); external_scope_metadata_.clear(); } + { + std::lock_guard lock(memory_allocation_mutex_); + memory_allocations_.clear(); + } { std::lock_guard lock(agent_mutex_); gpu_device_ids_.clear(); @@ -310,6 +317,19 @@ int RocprofilerBackend::toolInitialize() { &reason)) { return -1; } + if (opts_.enable_memory_tracking) { + const auto status = + rocprofiler_configure_buffer_tracing_service( + context_, ROCPROFILER_BUFFER_TRACING_MEMORY_ALLOCATION, + nullptr, 0, buffer_); + if (status == ROCPROFILER_STATUS_SUCCESS) { + memory_activity_configured_.store(true, std::memory_order_release); + } else { + GFL_LOG_WARN( + "[ROCProfilerBackend] memory-allocation tracing unavailable: ", + StatusToString(status)); + } + } // Resolve the user-facing request to an AMD-native path before creating an // engine. Unsupported requests remain trace-capable and are reported as @@ -518,6 +538,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 memory_activity_rows = + memory_activity_rows_emitted_.load(std::memory_order_relaxed); const uint64_t dropped_records = trace_records_dropped_.load(std::memory_order_relaxed); const uint64_t queue_dropped_records = @@ -536,10 +558,15 @@ void RocprofilerBackend::emitCapabilities() { input.plan = resolvedPlan(); input.trace_configured = tool_registered_.load(std::memory_order_acquire); + input.memory_activity_requested = opts_.enable_memory_tracking; + input.memory_activity_configured = + memory_activity_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.memory_activity_rows = delta( + memory_activity_rows, capability_memory_activity_rows_baseline_); input.profiling_sample_rows = engine_ && engine_->hasData() ? 1 : 0; input.pm_sample_rows = @@ -561,6 +588,7 @@ void RocprofilerBackend::emitCapabilities() { capability_kernel_rows_baseline_ = kernel_rows; capability_memcpy_rows_baseline_ = memcpy_rows; + capability_memory_activity_rows_baseline_ = memory_activity_rows; capability_pm_sample_rows_baseline_ = pm_sample_rows; capability_dropped_records_baseline_ = dropped_records; capability_queue_dropped_records_baseline_ = queue_dropped_records; @@ -766,6 +794,14 @@ void RocprofilerBackend::bufferTracingShim(rocprofiler_context_id_t, backend->handleMemoryCopy(*record); break; } + case ROCPROFILER_BUFFER_TRACING_MEMORY_ALLOCATION: { + const auto* record = static_cast< + const rocprofiler_buffer_tracing_memory_allocation_record_t*>( + header->payload); + backend->handleMemoryAllocation(*record); + break; + } + default: break; } @@ -1135,4 +1171,88 @@ void RocprofilerBackend::handleMemoryCopy( } } +void RocprofilerBackend::handleMemoryAllocation( + const rocprofiler_buffer_tracing_memory_allocation_record_t& data) { + const auto operation = NormalizeAmdMemoryAllocationOperation( + static_cast(data.operation)); + if (!operation.has_value()) return; + + const uint64_t address = data.address.value; + const auto endpoint = resolveTraceEndpoint(data.agent_id); + const auto direct_device_id = + ResolveAmdMemoryAllocationDeviceId(endpoint); + const bool direct_host = endpoint.kind == AmdTraceAgentKind::Cpu; + + MemoryAllocationMetadata metadata; + bool attributed = false; + if (*operation == 1) { + if (direct_device_id.has_value() || direct_host) { + metadata.device_id = direct_device_id.value_or(0); + metadata.memory_kind = ResolveAmdMemoryAllocationKind(endpoint); + metadata.bytes = data.allocation_size; + attributed = true; + if (address != 0) { + std::lock_guard lock(memory_allocation_mutex_); + memory_allocations_[address] = metadata; + } + } + } else { + if (address != 0) { + std::lock_guard lock(memory_allocation_mutex_); + if (const auto itr = memory_allocations_.find(address); + itr != memory_allocations_.end()) { + metadata = itr->second; + memory_allocations_.erase(itr); + attributed = true; + } + } + if (!attributed && (direct_device_id.has_value() || direct_host)) { + metadata.device_id = direct_device_id.value_or(0); + metadata.memory_kind = ResolveAmdMemoryAllocationKind(endpoint); + metadata.bytes = data.allocation_size; + attributed = true; + } + } + + if (!attributed) { + trace_records_unattributed_.fetch_add(1, std::memory_order_relaxed); + return; + } + + ActivityRecord out{}; + out.type = TraceType::MEMORY_ALLOC; + out.device_id = metadata.device_id; + out.cpu_start_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.memory_op = *operation; + out.memory_kind = metadata.memory_kind; + out.address = address; + out.bytes = data.allocation_size > 0 ? data.allocation_size : metadata.bytes; + 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)) { + memory_activity_rows_emitted_.fetch_add(1, + std::memory_order_relaxed); + } else { + trace_records_queue_dropped_.fetch_add(1, std::memory_order_relaxed); + } +} + } // namespace gpufl::amd diff --git a/include/gpufl/backends/amd/rocprofiler_backend.hpp b/include/gpufl/backends/amd/rocprofiler_backend.hpp index 9c9db66..0c5fb72 100644 --- a/include/gpufl/backends/amd/rocprofiler_backend.hpp +++ b/include/gpufl/backends/amd/rocprofiler_backend.hpp @@ -72,6 +72,12 @@ class RocprofilerBackend final : public IMonitorBackend { int scope_depth = 0; }; + struct MemoryAllocationMetadata { + uint32_t device_id = 0; + uint8_t memory_kind = 0; + uint64_t bytes = 0; + }; + bool configureRocprofiler(const MonitorOptions& opts, std::string* reason); void resetToolState(); bool registerTool(std::string* reason); @@ -86,6 +92,8 @@ 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 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; @@ -120,6 +128,8 @@ class RocprofilerBackend final : public IMonitorBackend { mutable std::unordered_map demangle_cache_; mutable std::mutex external_scope_mutex_; std::unordered_map external_scope_metadata_; + mutable std::mutex memory_allocation_mutex_; + std::unordered_map memory_allocations_; mutable std::mutex agent_mutex_; std::unordered_map gpu_device_ids_; @@ -154,6 +164,7 @@ class RocprofilerBackend final : public IMonitorBackend { std::atomic kernel_rows_emitted_{0}; std::atomic memcpy_rows_emitted_{0}; + std::atomic memory_activity_rows_emitted_{0}; std::atomic trace_records_dropped_{0}; std::atomic trace_records_queue_dropped_{0}; std::atomic trace_buffer_flush_failures_{0}; @@ -164,6 +175,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_memory_activity_rows_baseline_ = 0; mutable uint64_t capability_pm_sample_rows_baseline_ = 0; mutable uint64_t capability_dropped_records_baseline_ = 0; mutable uint64_t capability_queue_dropped_records_baseline_ = 0; @@ -179,6 +191,7 @@ class RocprofilerBackend final : public IMonitorBackend { std::atomic start_failure_logged_{false}; std::mutex start_stop_mutex_; std::atomic tool_registered_{false}; + std::atomic memory_activity_configured_{false}; }; } // namespace gpufl::amd diff --git a/include/gpufl/core/activity_record.hpp b/include/gpufl/core/activity_record.hpp index 831a477..e7ad19c 100644 --- a/include/gpufl/core/activity_record.hpp +++ b/include/gpufl/core/activity_record.hpp @@ -107,20 +107,18 @@ struct ActivityRecord { // Memory-allocation specific. // - // memory_op: 1 = ALLOC, 2 = FREE. Mirrors CUpti_ActivityMemoryOperationType - // (we collapse the variants we don't surface in v1 - release-async - // and so on - into the two top-level buckets; the dashboard only - // distinguishes alloc-vs-free in its first iteration). + // memory_op is the portable 1 = ALLOC, 2 = FREE contract. Backends + // collapse native variants into those two top-level operations. // - // memory_kind values mirror CUpti_ActivityMemoryKind: + // memory_kind retains CUPTI-compatible wire values so existing consumers + // remain cross-vendor: // 0 = UNKNOWN, 1 = PAGEABLE_HOST, 2 = PINNED_HOST, 3 = DEVICE, // 4 = ARRAY, 5 = MANAGED, 6 = DEVICE_STATIC, 7 = MANAGED_STATIC. - // Backend stores the raw integer; frontend maps to a label. + // ROCprofiler identifies only the owning agent, so AMD GPU allocations + // map to DEVICE and CPU allocations remain UNKNOWN. // - // address is the GPU virtual address (uint64) returned by cudaMalloc - // (or freed by cudaFree). Stored as uint64 because Blackwell-class - // addresses can be large; the dashboard renders it as 0x-prefixed - // hex. + // address is the uint64 virtual address allocated or freed. The + // dashboard renders it as 0x-prefixed hexadecimal text. uint8_t memory_op = 0; uint8_t memory_kind = 0; uint64_t address = 0; diff --git a/include/gpufl/core/events/memory_events.hpp b/include/gpufl/core/events/memory_events.hpp index 7c63671..02eca4b 100644 --- a/include/gpufl/core/events/memory_events.hpp +++ b/include/gpufl/core/events/memory_events.hpp @@ -51,16 +51,15 @@ struct MemsetEvent { uint64_t bytes = 0; }; -// One CUPTI MEMORY2 record - `cudaMalloc` / `cudaFree` / `cudaMallocAsync` / -// etc. Replaces per-event `memory_alloc_event` JSON with a packed row -// inside `memory_alloc_event_batch`. Pure-numeric fields → no dictionary -// encoding, just envelope amortization. Saves ~85% on alloc-heavy -// workloads. +// One CUPTI MEMORY2 or ROCprofiler memory-allocation record. Packed rows +// replace per-event `memory_alloc_event` JSON inside +// `memory_alloc_event_batch`. Pure-numeric fields need no dictionary +// encoding, which keeps allocation-heavy workloads compact. struct MemoryAllocEventBatchRow { int64_t start_ns = 0; - int64_t duration_ns = 0; // 0 in v1 - CUPTI doesn't emit alloc duration + int64_t duration_ns = 0; // host-call duration when available uint8_t memory_op = 0; // 1=ALLOC, 2=FREE - uint8_t memory_kind = 0; // CUpti_ActivityMemoryKind + uint8_t memory_kind = 0; // portable values; see ActivityRecord uint64_t address = 0; // GPU virtual address uint64_t bytes = 0; uint32_t device_id = 0; @@ -79,27 +78,17 @@ struct MemcpyBatchRow { }; /** - * One CUDA memory-management event captured by CUPTI's - * CUPTI_ACTIVITY_KIND_MEMORY2 stream. + * One GPU memory-management event captured by CUPTI or ROCprofiler. * - * Covers cudaMalloc / cudaFree / cudaMallocAsync / cudaFreeAsync / - * cudaMallocManaged / cudaMallocHost (and their driver-API cousins). - * One event per call. Note that cudaMallocAsync is associated with a - * stream and the reported {@code start_ns} is the host call time - * (not the GPU completion time) - the host-side cost is what users - * actually pay for in their python/c++ code. + * CUDA covers allocation variants such as cudaMalloc, cudaFree, and their + * asynchronous or managed-memory forms. AMD ROCprofiler reports ALLOCATE, + * VMEM_ALLOCATE, FREE, and VMEM_FREE; the backend normalizes those to the + * same portable alloc/free operation values. * - * Per-event JSON. Volume in PyTorch workloads is typically <1k events - * per session because torch's caching allocator absorbs most python- - * level allocations; only large-block CUDA-level mallocs reach this - * stream. TensorFlow eager mode is the high-volume edge case - if it - * becomes a problem the gating flag {@code enable_memory_tracking} - * lets users opt out without losing other CUPTI streams. - * - * The {@code address} field is the VA returned by cudaMalloc (or - * being freed by cudaFree). Pairing alloc → free across the session - * for leak / fragmentation analysis is a v2 follow-up; v1 just - * stores raw events. + * The reported timestamps describe the host call. The address and byte + * count allow alloc/free pairing and future leak or fragmentation analysis. + * Allocation tracking can be disabled with enable_memory_tracking without + * disabling the other activity streams. */ struct MemoryAllocEvent { int pid = 0; @@ -108,7 +97,7 @@ struct MemoryAllocEvent { int64_t start_ns = 0; int64_t duration_ns = 0; // host-side; usually tiny but non-zero uint8_t memory_op = 0; // 1 = ALLOC, 2 = FREE - uint8_t memory_kind = 0; // CUpti_ActivityMemoryKind + uint8_t memory_kind = 0; // portable CUPTI-compatible kind value uint64_t address = 0; uint64_t bytes = 0; uint32_t device_id = 0; diff --git a/include/gpufl/core/model/batch_models.cpp b/include/gpufl/core/model/batch_models.cpp index 9f14241..0dd655a 100644 --- a/include/gpufl/core/model/batch_models.cpp +++ b/include/gpufl/core/model/batch_models.cpp @@ -348,8 +348,8 @@ std::string MemoryAllocEventBatchModel::buildJson() const { if (rows.empty()) return {}; const int64_t base = rows.front().start_ns; - // All numeric fields → no dictionary lookup needed. address is a - // 64-bit GPU virtual address (can exceed 32 bits on Blackwell). + // All numeric fields → no dictionary lookup needed. Address is a + // 64-bit GPU virtual address and may exceed 32 bits. std::ostringstream oss; oss << "{\"version\":1,\"type\":\"memory_alloc_event_batch\"" << ",\"session_id\":\"" << jsonEscape(session_id_) << '"' diff --git a/include/gpufl/core/monitor.cpp b/include/gpufl/core/monitor.cpp index 5a1dff3..86c0f5a 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::MemoryAllocRowsSeen() { + return g_state.batches.memoryAllocRowsSeen(); +} + void Monitor::EmitPmSamplingConfig(uint32_t device_id, uint32_t interval_us, uint32_t max_samples, const std::string& preset, const std::vector& metrics) { const Runtime* rt = runtime(); const auto segment = rt ? rt->acquireSegmentContext() : nullptr; diff --git a/include/gpufl/core/monitor.hpp b/include/gpufl/core/monitor.hpp index 6f332fb..7e9cc4d 100644 --- a/include/gpufl/core/monitor.hpp +++ b/include/gpufl/core/monitor.hpp @@ -458,6 +458,8 @@ class Monitor { static uint64_t ScopeAttributionTruncated(); /** @brief PM metric rows that passed through scope attribution. */ static uint64_t PmSampleRowsSeen(); + /** @brief Memory-allocation rows accepted by the batch pipeline. */ + static uint64_t MemoryAllocRowsSeen(); /** * Flush all currently available records and batches before publishing a diff --git a/include/gpufl/core/monitor_batch_manager.cpp b/include/gpufl/core/monitor_batch_manager.cpp index 973ba21..fafb3c2 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(); + memoryAllocRowsSeen_.store(0, std::memory_order_relaxed); pendingDetails_.clear(); kernelBatchId_ = 0; @@ -432,6 +433,10 @@ uint64_t MonitorBatchManager::pmSampleRowsSeen() const { return pmSampleRowsSeen_; } +uint64_t MonitorBatchManager::memoryAllocRowsSeen() const { + return memoryAllocRowsSeen_.load(std::memory_order_relaxed); +} + void MonitorBatchManager::resolveScopeIdsForTesting(std::vector& rows, uint32_t fallback_id) { std::vector candidates; @@ -604,6 +609,7 @@ void MonitorBatchManager::resolveScopeIdsForBatch(std::vector& cand } bool MonitorBatchManager::pushMemoryAlloc(const MemoryAllocEventBatchRow& row) { + memoryAllocRowsSeen_.fetch_add(1, std::memory_order_relaxed); memAllocBatch_.push(row); return memAllocBatch_.needsFlush(); } diff --git a/include/gpufl/core/monitor_batch_manager.hpp b/include/gpufl/core/monitor_batch_manager.hpp index 3762205..4d0e4f3 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 Memory-allocation rows accepted by the batch pipeline. */ + uint64_t memoryAllocRowsSeen() const; /** @brief Test seam: resolve a batch exactly as the drain path does. */ void resolveScopeIdsForTesting(std::vector& rows, uint32_t fallback_id); @@ -225,6 +227,7 @@ class MonitorBatchManager { uint64_t pmSampleRowsSeen_ = 0; BatchBuffer syncBatch_; + std::atomic memoryAllocRowsSeen_{0}; BatchBuffer memAllocBatch_; uint64_t syncBatchId_ = 0; uint64_t memAllocBatchId_ = 0; diff --git a/include/gpufl/core/trace_type.hpp b/include/gpufl/core/trace_type.hpp index 66daa44..1a3291f 100644 --- a/include/gpufl/core/trace_type.hpp +++ b/include/gpufl/core/trace_type.hpp @@ -42,26 +42,20 @@ enum class TraceType : uint8_t { // event-wait); see ActivityRecord::sync_type // sync_event_id - for event-based syncs; 0 for stream/device SYNCHRONIZATION, - // CUDA memory allocation event captured via - // CUPTI_ACTIVITY_KIND_MEMORY2. One record per cudaMalloc / cudaFree - // / cudaMallocAsync / cudaFreeAsync / cudaMallocManaged / cudaMallocHost. - // Volume per session is mid-low (typically <1k events even for big - // workloads - PyTorch's caching allocator absorbs most fine-grained - // allocations into a few large CUDA-level blocks), so we emit per-event - // JSON (no batching). + // GPU memory-allocation event captured via CUPTI MEMORY2 or + // ROCprofiler memory-allocation tracing. Rows are normalized and batched + // into the shared memory_alloc_event_batch wire format. // // Fields used on ActivityRecord: // cpu_start_ns - host call wall time + // duration_ns - host call duration when the backend supplies it // bytes - size of the allocation - // device_id - target device (0 for host allocs) - // corr_id - CUPTI correlationId, useful for joining to the - // API call that triggered the alloc - // stream - stream id for cudaMallocAsync; 0 otherwise - // memory_op - alloc / free (uint8_t enum) - // memory_kind - device / host / managed / pinned (uint8_t enum) - // address - VA address of the allocation; for free, the - // address being freed (lets us pair alloc/free - // in a future v2 backend pass) + // device_id - target device (0 for host allocations) + // corr_id - backend correlation id + // stream - stream id for asynchronous CUDA allocation; 0 on AMD + // memory_op - portable 1=ALLOC, 2=FREE value + // memory_kind - portable device / host / managed / pinned value + // address - VA allocated or freed MEMORY_ALLOC, // F4: CUDA graph launch event captured via // CUPTI_ACTIVITY_KIND_GRAPH_TRACE. One record per cudaGraphLaunch diff --git a/include/gpufl/report/text_report.cpp b/include/gpufl/report/text_report.cpp index 47b0dd0..565e46d 100644 --- a/include/gpufl/report/text_report.cpp +++ b/include/gpufl/report/text_report.cpp @@ -164,6 +164,20 @@ std::string resolveCopyKind(int kind) { return (it != kCopyKindNames.end()) ? it->second : "Unknown(" + std::to_string(kind) + ")"; } +std::string resolveMemoryKind(int kind) { + switch (kind) { + case 0: return "Unknown"; + case 1: return "Pageable"; + case 2: return "Pinned"; + case 3: return "Device"; + case 4: return "Array"; + case 5: return "Managed"; + case 6: return "DeviceStatic"; + case 7: return "ManagedStatic"; + default: return "Unknown(" + std::to_string(kind) + ")"; + } +} + // Values match CUpti_ActivityPCSamplingStallReason enum from cupti_activity.h const std::map kStallNames = { {2, "Instruction Fetch"}, {3, "Execution Dependency"}, @@ -397,12 +411,43 @@ bool TextReport::isAmdSession() const { identity.find("advanced micro devices") != std::string::npos; } +bool TextReport::tryParseMemoryAllocationRecord(const JsonValue& rec) { + const std::string type = rec.value("type", ""); + if (type == "memory_alloc_event_batch") { + auto ci = buildColumnIndex(rec["columns"]); + int64_t base = rec.value("base_time_ns", 0); + for (const auto& row : rec["rows"].get_array()) { + memory_allocations_.push_back({ + base + rowInt(row, ci, "dt_ns"), + static_cast(rowInt(row, ci, "memory_op")), + static_cast(rowInt(row, ci, "memory_kind")), + rowU64(row, ci, "address"), + rowU64(row, ci, "bytes"), + }); + } + return true; + } + + if (type == "memory_alloc_event") { + memory_allocations_.push_back({ + rec.value("start_ns", 0), + static_cast(rec.value("memory_op", 0)), + static_cast(rec.value("memory_kind", 0)), + rec.value("address", 0), + rec.value("bytes", 0), + }); + return true; + } + return false; +} + void TextReport::parseDeviceLog(const std::vector& records, const std::unordered_map& kernel_dict) { std::unordered_map details; for (const auto& rec : records) { const std::string type = rec.value("type", ""); + if (tryParseMemoryAllocationRecord(rec)) continue; if ((type == "job_start" || type == "init") && info_.app_name.empty()) { parseJobStart(rec, info_); @@ -506,6 +551,7 @@ void TextReport::parseScopeLog(const std::vector& records, for (const auto& rec : records) { const std::string type = rec.value("type", ""); + if (tryParseMemoryAllocationRecord(rec)) continue; if ((type == "job_start" || type == "init") && info_.app_name.empty()) { parseJobStart(rec, info_); @@ -641,6 +687,7 @@ std::string TextReport::generate() const { writeTopKernels(out); writeKernelDetails(out); writeMemcpySummary(out); + writeMemoryAllocationSummary(out); writeSystemMetrics(out); writeScopeSummary(out); writePerfMetricsSummary(out); @@ -916,6 +963,96 @@ void TextReport::writeMemcpySummary(std::ostringstream& out) const { } } +void TextReport::writeMemoryAllocationSummary(std::ostringstream& out) const { + out << "\n" << SEP << "\n Memory Allocation Summary\n" << SEP << "\n"; + if (memory_allocations_.empty()) { + out << " (No memory allocation data)\n"; + return; + } + + struct KindStats { + size_t allocations = 0; + size_t frees = 0; + uint64_t bytes_allocated = 0; + }; + + std::vector ordered; + ordered.reserve(memory_allocations_.size()); + for (const auto& record : memory_allocations_) ordered.push_back(&record); + std::stable_sort(ordered.begin(), ordered.end(), + [](const auto* lhs, const auto* rhs) { + return lhs->start_ns < rhs->start_ns; + }); + + std::map grouped; + std::unordered_map live_bytes; + uint64_t bytes_allocated = 0; + uint64_t bytes_freed = 0; + uint64_t current_live_bytes = 0; + uint64_t peak_live_bytes = 0; + size_t allocation_count = 0; + size_t free_count = 0; + + for (const auto* record : ordered) { + if (record->memory_op == 1) { + ++allocation_count; + auto& stats = grouped[record->memory_kind]; + ++stats.allocations; + stats.bytes_allocated += record->bytes; + bytes_allocated += record->bytes; + + if (record->address != 0) { + auto existing = live_bytes.find(record->address); + if (existing != live_bytes.end()) { + current_live_bytes -= + (std::min)(current_live_bytes, existing->second); + } + live_bytes[record->address] = record->bytes; + current_live_bytes += record->bytes; + peak_live_bytes = (std::max)(peak_live_bytes, current_live_bytes); + } + } else if (record->memory_op == 2) { + ++free_count; + ++grouped[record->memory_kind].frees; + + uint64_t freed_bytes = record->bytes; + if (record->address != 0) { + auto existing = live_bytes.find(record->address); + if (existing != live_bytes.end()) { + if (freed_bytes == 0) freed_bytes = existing->second; + current_live_bytes -= + (std::min)(current_live_bytes, existing->second); + live_bytes.erase(existing); + } + } + bytes_freed += freed_bytes; + } + } + + out << " Total Events: " << memory_allocations_.size() << "\n"; + out << " Allocations: " << allocation_count << "\n"; + out << " Frees: " << free_count << "\n"; + out << " Bytes Allocated: " << fmtBytes(bytes_allocated) << "\n"; + out << " Bytes Freed: " << fmtBytes(bytes_freed) << "\n"; + out << " Peak Tracked Live: " << fmtBytes(peak_live_bytes) << "\n\n"; + + out << " By Memory Kind:\n"; + out << " " << std::left << std::setw(18) << "Kind" + << std::right << std::setw(13) << "Allocations" + << std::setw(10) << "Frees" + << std::setw(18) << "Bytes Allocated" << "\n"; + out << " " << std::string(59, '-') << "\n"; + for (const auto& [kind, stats] : grouped) { + out << " " << std::left << std::setw(18) << resolveMemoryKind(kind) + << std::right << std::setw(13) << stats.allocations + << std::setw(10) << stats.frees + << std::setw(18) << fmtBytes(stats.bytes_allocated) << "\n"; + } + + out << "\n Note: Peak tracked live memory is derived from matched observed addresses.\n" + << " Runtime and allocator-internal activity may be included.\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 aa5b1c2..ee3fca0 100644 --- a/include/gpufl/report/text_report.hpp +++ b/include/gpufl/report/text_report.hpp @@ -63,6 +63,14 @@ class TextReport { int copy_kind = 0; }; + struct MemoryAllocationRecord { + int64_t start_ns = 0; + uint8_t memory_op = 0; + uint8_t memory_kind = 0; + uint64_t address = 0; + uint64_t bytes = 0; + }; + struct DeviceMetricRecord { int64_t ts_ns = 0; int gpu_util = 0; @@ -171,6 +179,7 @@ class TextReport { SessionInfo info_; std::vector kernels_; std::vector memcpy_; + std::vector memory_allocations_; std::vector device_metrics_; std::vector host_metrics_; std::vector scope_events_; @@ -193,6 +202,7 @@ class TextReport { std::unordered_map& metric_dict); void parseDeviceLog(const std::vector& records, const std::unordered_map& kernel_dict); + bool tryParseMemoryAllocationRecord(const JsonValue& record); void parseScopeLog(const std::vector& records, const std::unordered_map& scope_name_dict, const std::unordered_map& function_dict, @@ -212,6 +222,7 @@ class TextReport { void writeTopKernels(std::ostringstream& out) const; void writeKernelDetails(std::ostringstream& out) const; void writeMemcpySummary(std::ostringstream& out) const; + void writeMemoryAllocationSummary(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 744d67f..36023c0 100644 --- a/tests/backends/amd/test_amd_profiling_policy.cpp +++ b/tests/backends/amd/test_amd_profiling_policy.cpp @@ -301,3 +301,42 @@ TEST(AmdCaptureCapabilities, ClientQueueDropsDegradeEndToEndDelivery) { ASSERT_NE(correlation, nullptr); EXPECT_EQ(correlation->status, "enabled"); } + +TEST(AmdCaptureCapabilities, MemoryActivityReportsCollectionAndAvailability) { + gpufl::amd::AmdCaptureCapabilityInput input; + input.session_id = "amd-session"; + input.plan = gpufl::amd::ResolveAmdProfilingPlan( + gpufl::ProfilingEngine::Trace, {}); + input.trace_configured = true; + input.memory_activity_requested = true; + input.memory_activity_configured = true; + input.memory_activity_rows = 4; + + auto event = gpufl::amd::BuildAmdCaptureCapabilitiesEvent(input); + const auto* memory = FindCapability(event, "memory_activity"); + ASSERT_NE(memory, nullptr); + EXPECT_TRUE(memory->requested); + EXPECT_EQ(memory->status, "collected"); + EXPECT_EQ(memory->mode, "rocprofiler_memory_allocation"); + + input.memory_activity_rows = 0; + event = gpufl::amd::BuildAmdCaptureCapabilitiesEvent(input); + memory = FindCapability(event, "memory_activity"); + ASSERT_NE(memory, nullptr); + EXPECT_EQ(memory->status, "enabled_no_data"); + EXPECT_EQ(memory->reason_code, "enabled_but_no_records"); + + input.memory_activity_configured = false; + event = gpufl::amd::BuildAmdCaptureCapabilitiesEvent(input); + memory = FindCapability(event, "memory_activity"); + ASSERT_NE(memory, nullptr); + EXPECT_EQ(memory->status, "skipped"); + EXPECT_EQ(memory->reason_code, + "rocprofiler_memory_allocation_unavailable"); + + input.memory_activity_requested = false; + event = gpufl::amd::BuildAmdCaptureCapabilitiesEvent(input); + memory = FindCapability(event, "memory_activity"); + ASSERT_NE(memory, nullptr); + EXPECT_EQ(memory->status, "not_requested"); +} diff --git a/tests/backends/amd/test_amd_trace_policy.cpp b/tests/backends/amd/test_amd_trace_policy.cpp index a99adac..00dabba 100644 --- a/tests/backends/amd/test_amd_trace_policy.cpp +++ b/tests/backends/amd/test_amd_trace_policy.cpp @@ -39,3 +39,31 @@ TEST(AmdTracePolicy, HostOnlyAndUnmappedGpuCopiesAreUnattributed) { EXPECT_FALSE(gpufl::amd::ResolveAmdMemoryCopyDeviceId(unmapped_gpu, Cpu()) .has_value()); } + +TEST(AmdTracePolicy, MemoryAllocationUsesOwningGpuDevice) { + EXPECT_EQ(gpufl::amd::ResolveAmdMemoryAllocationDeviceId(Gpu(6)), 6u); + EXPECT_FALSE( + gpufl::amd::ResolveAmdMemoryAllocationDeviceId(Cpu()).has_value()); + + const gpufl::amd::AmdTraceEndpoint unmapped_gpu{ + gpufl::amd::AmdTraceAgentKind::Gpu, std::nullopt}; + EXPECT_FALSE(gpufl::amd::ResolveAmdMemoryAllocationDeviceId(unmapped_gpu) + .has_value()); +} + +TEST(AmdTracePolicy, MemoryAllocationOperationsUsePortableWireValues) { + EXPECT_EQ(gpufl::amd::NormalizeAmdMemoryAllocationOperation(1), 1u); + EXPECT_EQ(gpufl::amd::NormalizeAmdMemoryAllocationOperation(2), 1u); + EXPECT_EQ(gpufl::amd::NormalizeAmdMemoryAllocationOperation(3), 2u); + EXPECT_EQ(gpufl::amd::NormalizeAmdMemoryAllocationOperation(4), 2u); + EXPECT_FALSE( + gpufl::amd::NormalizeAmdMemoryAllocationOperation(0).has_value()); + EXPECT_FALSE( + gpufl::amd::NormalizeAmdMemoryAllocationOperation(5).has_value()); +} + +TEST(AmdTracePolicy, MemoryAllocationKindIsTruthfulAcrossAgents) { + EXPECT_EQ(gpufl::amd::ResolveAmdMemoryAllocationKind(Gpu(0)), 3u); + EXPECT_EQ(gpufl::amd::ResolveAmdMemoryAllocationKind(Cpu()), 0u); + EXPECT_EQ(gpufl::amd::ResolveAmdMemoryAllocationKind({}), 0u); +} diff --git a/tests/core/test_monitor.cpp b/tests/core/test_monitor.cpp index 68865e9..5ed22f8 100644 --- a/tests/core/test_monitor.cpp +++ b/tests/core/test_monitor.cpp @@ -505,3 +505,21 @@ TEST(ScopeAttributionTest, OldTraceHistoryEvictedDuringPmIsNotPartialAttribution EXPECT_EQ(m.scopeAttributionTruncated(), 0u); } + +TEST(MemoryAllocationBatchTest, CountsAcceptedRowsAndResets) { + gpufl::detail::MonitorBatchManager manager; + manager.reset(); + + EXPECT_EQ(manager.memoryAllocRowsSeen(), 0u); + gpufl::MemoryAllocEventBatchRow row{}; + row.start_ns = 100; + (void)manager.pushMemoryAlloc(row); + EXPECT_EQ(manager.memoryAllocRowsSeen(), 1u); + + row.start_ns = 200; + (void)manager.pushMemoryAlloc(row); + EXPECT_EQ(manager.memoryAllocRowsSeen(), 2u); + + manager.reset(); + EXPECT_EQ(manager.memoryAllocRowsSeen(), 0u); +} diff --git a/tests/core/test_text_report.cpp b/tests/core/test_text_report.cpp index 2e50353..97574c4 100644 --- a/tests/core/test_text_report.cpp +++ b/tests/core/test_text_report.cpp @@ -93,4 +93,43 @@ TEST_F(TextReportTest, MeaningfulProfileRowsUseBackendNeutralSectionNames) { EXPECT_EQ(report.find("Other SASS Metrics:"), std::string::npos); } +TEST_F(TextReportTest, MemoryAllocationBatchRowsProduceSummary) { + WriteLog("scope", { + R"({"type":"job_start","session_id":"s1","app":"memory_report","ts_ns":1000,"gpu_static_devices":[{"name":"AMD Radeon Test","vendor":"AMD","multi_processor_count":32}]})", + R"({"type":"memory_alloc_event_batch","session_id":"s1","base_time_ns":1000,"columns":["dt_ns","duration_ns","memory_op","memory_kind","address","bytes","device_id","stream_id","corr_id"],"rows":[[10,1,1,3,4096,1024,0,0,1],[20,1,1,3,8192,2048,0,0,2],[30,1,2,3,4096,1024,0,0,3],[40,1,1,0,12288,512,0,0,4],[50,1,2,3,8192,2048,0,0,5]]})", + R"({"type":"shutdown","session_id":"s1","ts_ns":3000})", + }); + + const std::string report = Generate(); + EXPECT_NE(report.find("Memory Allocation Summary"), std::string::npos); + EXPECT_NE(report.find("Total Events: 5"), std::string::npos); + EXPECT_NE(report.find("Allocations: 3"), std::string::npos); + EXPECT_NE(report.find("Frees: 2"), std::string::npos); + EXPECT_NE(report.find("Bytes Allocated: 3.5 KB"), std::string::npos); + EXPECT_NE(report.find("Bytes Freed: 3.0 KB"), std::string::npos); + EXPECT_NE(report.find("Peak Tracked Live: 3.0 KB"), std::string::npos); + EXPECT_NE(report.find("By Memory Kind:"), std::string::npos); + EXPECT_NE(report.find("Device"), std::string::npos); + EXPECT_NE(report.find("Unknown"), std::string::npos); + + const auto allocations = report.find("Memory Allocation Summary"); + const auto system_metrics = report.find("System Metrics"); + EXPECT_LT(allocations, system_metrics); +} + +TEST_F(TextReportTest, LegacyMemoryAllocationEventProducesSummary) { + WriteLog("device", { + R"({"type":"job_start","session_id":"s1","app":"legacy_memory_report","ts_ns":1000,"gpu_static_devices":[{"name":"NVIDIA Test GPU","vendor":"NVIDIA","multi_processor_count":10}]})", + R"({"type":"memory_alloc_event","session_id":"s1","start_ns":1100,"duration_ns":0,"memory_op":1,"memory_kind":5,"address":4096,"bytes":2048,"device_id":0,"stream_id":0,"corr_id":1})", + R"({"type":"shutdown","session_id":"s1","ts_ns":3000})", + }); + + const std::string report = Generate(); + EXPECT_NE(report.find("Total Events: 1"), std::string::npos); + EXPECT_NE(report.find("Allocations: 1"), std::string::npos); + EXPECT_NE(report.find("Frees: 0"), std::string::npos); + EXPECT_NE(report.find("Bytes Allocated: 2.0 KB"), std::string::npos); + EXPECT_NE(report.find("Managed"), std::string::npos); +} + } // namespace