feat(amd): add memory allocation tracing

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)

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

example/amd/memory_allocation_rows.cpp

@@ -0,0 +1,152 @@
+#include <hip/hip_runtime.h>
+
+#include <chrono>
+#include <cstdint>
+#include <iostream>
+#include <thread>
+#include <vector>
+
+#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<void*> 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<size_t>(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;
+}

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

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;

include/gpufl/backends/amd/amd_trace_policy.cpp

@@ -21,4 +21,23 @@ std::optional<uint32_t> ResolveAmdMemoryCopyDeviceId(
     return std::nullopt;
 }
 
+std::optional<uint32_t> ResolveAmdMemoryAllocationDeviceId(
+    const AmdTraceEndpoint& agent) {
+    if (agent.kind != AmdTraceAgentKind::Gpu) return std::nullopt;
+    return agent.device_id;
+}
+
+std::optional<uint8_t> 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

include/gpufl/backends/amd/amd_trace_policy.hpp

@@ -26,4 +26,21 @@ std::optional<uint32_t> 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<uint32_t> 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<uint8_t> 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