diff --git a/CMakeLists.txt b/CMakeLists.txt index 6a1e2af..e47889d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -708,6 +708,7 @@ if(GPUFL_ENABLE_AMD) target_sources(gpufl PRIVATE include/gpufl/backends/amd/monitor_adapter_amd.cpp include/gpufl/backends/amd/rocprofiler_backend.cpp + include/gpufl/backends/amd/engine/device_counter_engine.cpp include/gpufl/backends/amd/engine/dispatch_counter_engine.cpp ) message(STATUS "Found ROCprofiler-SDK support") diff --git a/example/amd/CMakeLists.txt b/example/amd/CMakeLists.txt index 8a9b631..f8d0d49 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 + pm_sampling_sample_rows.cpp vector_add_benchmark.cpp ) @@ -58,14 +59,17 @@ 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_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_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_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 ce2599d..4a5db05 100644 --- a/example/amd/README.md +++ b/example/amd/README.md @@ -8,24 +8,28 @@ 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` +- Per-dispatch AMD hardware counters via ROCprofiler dispatch counting +- Device-wide `PmSampling` timelines via ROCprofiler device counting - `gpufl` initialization with `backend = gpufl::BackendKind::Amd` - User-defined scope logging via `GFL_SCOPE(...)` - HIP example programs that run on ROCm hardware ## What Does Not Work Yet -- AMD profiling engines equivalent to CUPTI PC Sampling / SASS Metrics / Range Profiler +- AMD PC sampling +- Instruction-level SASS metrics and NVIDIA-compatible Range Profiler metrics Today, the AMD backend is useful for: - system metric logging - device inventory - automatic HIP kernel and memcpy tracing +- per-dispatch and device-wide hardware-counter profiling - scope-level application instrumentation It is not yet useful for: -- instruction-level or hardware-counter profiling +- PC sampling or instruction-level profiling ## Targets @@ -35,6 +39,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_pm_sampling_sample_rows` + - Selects AMD device counting and exits successfully only when each of two named scopes emits PM sample rows ## Build @@ -51,6 +57,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_pm_sampling_sample_rows ``` The AMD example targets are only added when CMake detects HIP successfully. @@ -106,8 +113,27 @@ 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_pm_sampling_sample_rows ``` +`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 +counting is unavailable or either scope does not produce a row. Its generated +report shows the same rows grouped by scope for manual inspection. + +The scope demo selects per-dispatch counters by default. To exercise the +device-wide PM timeline instead: + +```bash +GPUFL_PROFILING_ENGINE=PmSampling \ + ./build-rocm-examples/example/amd/amd_gpufl_scope_demo +``` + +`PmSampling` uses the portable `GPUBusy` derived counter by default. Set +`pm_sampling_metrics` programmatically to request other native ROCprofiler +counter names. + On a working ROCm system, `amd_check_device` should print output similar to: ```text @@ -123,6 +149,12 @@ Success! Device 0: AMD Radeon RX 9070 XT (arch gfx1201, capability 12.0) gfl_amd_scope ``` +`amd_pm_sampling_sample_rows` writes logs with prefix: + +```bash +gfl_amd_pm_rows +``` + With `rocprofiler-sdk` available, expect: - `job_start` inventory @@ -130,10 +162,12 @@ With `rocprofiler-sdk` available, expect: - `kernel_event_batch` - `kernel_detail` - `memcpy_event_batch` +- `profile_sample_batch` for dispatch-counting requests +- `pm_sampling_config` and `pm_sample_batch` for `PmSampling` - system metric samples - scope events Without `rocprofiler-sdk`, expect only telemetry, static inventory, and scope events. -Do not expect AMD profiling samples such as PC sampling or hardware counters. +PC samples and instruction-level SASS samples are not available on AMD yet. diff --git a/example/amd/pm_sampling_sample_rows.cpp b/example/amd/pm_sampling_sample_rows.cpp new file mode 100644 index 0000000..6501dc9 --- /dev/null +++ b/example/amd/pm_sampling_sample_rows.cpp @@ -0,0 +1,145 @@ +#include + +#include +#include + +#include "gpufl/core/monitor.hpp" +#include "gpufl/gpufl.hpp" + +namespace { + +constexpr int kElementCount = 1 << 20; +constexpr int kBlockSize = 256; +constexpr int kLaunchesPerPhase = 8; +constexpr int kIterationsPerLaunch = 1024; + +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 sampleRowsWorkload(float* values, const int count, + const int iterations) { + const int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index >= count) return; + + float x = values[index] + + static_cast((index & 1023) + 1) * 0.0001f; + float y = static_cast((threadIdx.x & 31) + 1) * 0.00001f; + for (int iteration = 0; iteration < iterations; ++iteration) { + x = x * 1.0000001f + y; + y = y * 0.9999999f + x * 0.000001f; + if (x > 4096.0f) x -= 4096.0f; + } + values[index] = x + y; +} + +bool RunPhase(float* values, const int launches, const int iterations) { + const dim3 block(kBlockSize); + const dim3 grid((kElementCount + block.x - 1) / block.x); + for (int launch = 0; launch < launches; ++launch) { + hipLaunchKernelGGL(sampleRowsWorkload, grid, block, 0, 0, values, + kElementCount, iterations); + } + return CheckHip(hipGetLastError(), "sampleRowsWorkload launch") && + CheckHip(hipDeviceSynchronize(), "sampleRowsWorkload sync"); +} + +} // namespace + +int main() { + gpufl::InitOptions opts; + opts.app_name = "amd_pm_sampling_sample_rows"; + opts.log_path = "gfl_amd_pm_rows"; + opts.backend = gpufl::BackendKind::Amd; + opts.profiling_engine = gpufl::ProfilingEngine::PmSampling; + opts.pm_sampling_interval_us = 1000; + opts.pm_sampling_max_samples = 4096; + opts.pm_sampling_preset = "overview"; + opts.pm_sampling_metrics = {"GPUBusy"}; + opts.pm_sampling_scope_only = 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 PM sampling\n"; + return 1; + } + + const std::string engine = + gpufl::Monitor::ResolvedProfilingEngineWireName(); + const bool engine_ok = engine == "amd.device_counting"; + std::cout << "=== GPUFL AMD PM Sample Rows ===\n" + << "Resolved engine: " << engine << "\n"; + if (!engine_ok) { + std::cerr << "Expected amd.device_counting; ROCprofiler device " + "counting may be unavailable\n"; + } + + float* device_values = 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)"); + } + + // Warm up HIP and load the kernel before opening a measured scope. This + // also gives ROCprofiler's deferred device-counting callback time to + // accept the configured profile. + if (workload_ok) workload_ok = RunPhase(device_values, 1, 64); + + const uint64_t rows_before = gpufl::Monitor::PmSampleRowsSeen(); + bool phase_a_ok = false; + if (workload_ok) { + GFL_SCOPE("pm_rows_phase_a") { + phase_a_ok = RunPhase(device_values, kLaunchesPerPhase, + kIterationsPerLaunch); + } + } + const uint64_t rows_after_a = gpufl::Monitor::PmSampleRowsSeen(); + + bool phase_b_ok = false; + if (workload_ok && phase_a_ok) { + GFL_SCOPE("pm_rows_phase_b") { + phase_b_ok = RunPhase(device_values, kLaunchesPerPhase, + kIterationsPerLaunch); + } + } + const uint64_t rows_after_b = gpufl::Monitor::PmSampleRowsSeen(); + + std::cout << "PM 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; + const bool phase_b_rows = rows_after_b > rows_after_a; + if (!phase_a_rows) { + std::cerr << "Phase A did not emit a PM sample row\n"; + } + if (!phase_b_rows) { + std::cerr << "Phase B did not emit a PM sample row\n"; + } + + if (device_values != nullptr) { + (void) hipFree(device_values); + } + + gpufl::shutdown(); + gpufl::generateReport(); + + const bool passed = engine_ok && workload_ok && phase_a_ok && phase_b_ok && + phase_a_rows && phase_b_rows; + if (!passed) return 2; + + std::cout << "\nPASS: both named scopes emitted GPUBusy sample rows.\n" + << "Inspect logs with prefix " << opts.log_path + << " for pm_sampling_config and pm_sample_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 fe7f8ae..a224a99 100644 --- a/include/gpufl/backends/amd/amd_capture_capabilities.cpp +++ b/include/gpufl/backends/amd/amd_capture_capabilities.cpp @@ -129,17 +129,20 @@ CaptureCapabilitiesEvent BuildAmdCaptureCapabilitiesEvent( !device_counting_requested ? "not_requested" : (device_counting_selected - ? (input.profiling_sample_rows > 0 ? "collected" - : "enabled_no_data") + ? (input.pm_sample_rows > 0 ? "collected" + : "enabled_no_data") : "skipped"), device_counting_selected ? "rocprofiler_device_counting_service" : "disabled", device_counting_selected - ? (input.profiling_sample_rows > 0 ? "" : "enabled_but_no_records") + ? (input.pm_sample_rows > 0 ? "" : "enabled_but_no_records") : (device_counting_requested ? input.plan.reason_code : "not_selected"), device_counting_selected - ? "AMD device counting was selected." + ? (input.pm_sample_rows > 0 + ? "Device-wide AMD hardware-counter samples were collected." + : "AMD device counting was selected but emitted no samples " + "this session.") : "AMD device counting is not available in the current " "implementation."); diff --git a/include/gpufl/backends/amd/amd_capture_capabilities.hpp b/include/gpufl/backends/amd/amd_capture_capabilities.hpp index 3f0e80d..cce4b01 100644 --- a/include/gpufl/backends/amd/amd_capture_capabilities.hpp +++ b/include/gpufl/backends/amd/amd_capture_capabilities.hpp @@ -15,6 +15,7 @@ struct AmdCaptureCapabilityInput { uint64_t kernel_rows = 0; uint64_t memcpy_rows = 0; uint64_t profiling_sample_rows = 0; + uint64_t pm_sample_rows = 0; uint64_t dropped_trace_records = 0; uint64_t dropped_client_records = 0; uint64_t trace_buffer_flush_failures = 0; diff --git a/include/gpufl/backends/amd/amd_profiling_policy.cpp b/include/gpufl/backends/amd/amd_profiling_policy.cpp index 8f0a32a..0f30d0f 100644 --- a/include/gpufl/backends/amd/amd_profiling_policy.cpp +++ b/include/gpufl/backends/amd/amd_profiling_policy.cpp @@ -1,5 +1,7 @@ #include "gpufl/backends/amd/amd_profiling_policy.hpp" +#include + namespace gpufl::amd { namespace { @@ -60,6 +62,22 @@ bool AmdRequestNeedsDeviceCounting(const ProfilingEngine engine) { engine == ProfilingEngine::Deep; } +std::vector ResolveAmdDeviceCountingMetrics( + const std::vector& requested_metrics) { + if (requested_metrics.empty()) return {"GPUBusy"}; + + std::vector resolved; + resolved.reserve(requested_metrics.size()); + std::unordered_set seen; + for (const auto& metric : requested_metrics) { + if (!metric.empty() && seen.emplace(metric).second) { + resolved.push_back(metric); + } + } + return resolved.empty() ? std::vector{"GPUBusy"} + : resolved; +} + std::optional ResolveAmdDispatchDeviceId( const uint64_t configured_agent_handle, const uint32_t configured_device_id, diff --git a/include/gpufl/backends/amd/amd_profiling_policy.hpp b/include/gpufl/backends/amd/amd_profiling_policy.hpp index dd8d321..58bdb7c 100644 --- a/include/gpufl/backends/amd/amd_profiling_policy.hpp +++ b/include/gpufl/backends/amd/amd_profiling_policy.hpp @@ -3,6 +3,7 @@ #include #include #include +#include #include "gpufl/core/monitor.hpp" @@ -43,6 +44,12 @@ bool AmdRequestNeedsDispatchCounting(ProfilingEngine engine); bool AmdRequestNeedsPcSampling(ProfilingEngine engine); bool AmdRequestNeedsDeviceCounting(ProfilingEngine engine); +// AMD's device-counting service consumes native ROCprofiler counter names. +// Explicit user metrics win; the default is a portable device-utilization +// signal that is available as a derived counter across supported AMD GPUs. +std::vector ResolveAmdDeviceCountingMetrics( + const std::vector& requested_metrics); + // Dispatch counting is currently configured for one GPU agent. Resolve only // records from that agent so a secondary GPU can never be mislabeled as device // zero (or as the configured primary device). diff --git a/include/gpufl/backends/amd/engine/amd_profiling_engine.hpp b/include/gpufl/backends/amd/engine/amd_profiling_engine.hpp index ea6bcbe..6997326 100644 --- a/include/gpufl/backends/amd/engine/amd_profiling_engine.hpp +++ b/include/gpufl/backends/amd/engine/amd_profiling_engine.hpp @@ -32,6 +32,10 @@ class AmdProfilingEngine { /// Periodically drain buffered profiling data into the monitor ring buffer. virtual void drain() = 0; + /// Cheap collector-loop service tick. Pull-based engines use this to + /// honor their sampling cadence; callback/buffer engines leave it idle. + virtual void service() {} + /// Release resources. virtual void shutdown() = 0; diff --git a/include/gpufl/backends/amd/engine/device_counter_engine.cpp b/include/gpufl/backends/amd/engine/device_counter_engine.cpp new file mode 100644 index 0000000..b7ad7ea --- /dev/null +++ b/include/gpufl/backends/amd/engine/device_counter_engine.cpp @@ -0,0 +1,412 @@ +#if !(GPUFL_ENABLE_AMD && GPUFL_HAS_ROCPROFILER_SDK) +#error "device_counter_engine.cpp requires GPUFL_ENABLE_AMD && GPUFL_HAS_ROCPROFILER_SDK" +#endif + +#include "gpufl/backends/amd/engine/device_counter_engine.hpp" + +#include +#include +#include + +#include "gpufl/backends/amd/amd_profiling_policy.hpp" +#include "gpufl/core/common.hpp" +#include "gpufl/core/debug_logger.hpp" +#include "gpufl/core/monitor.hpp" + +namespace gpufl::amd { +namespace { + +bool CheckStatus(const rocprofiler_status_t status, const char* call) { + if (status == ROCPROFILER_STATUS_SUCCESS) return true; + GFL_LOG_ERROR("[DeviceCounterEngine] ", call, " failed: status=", + static_cast(status)); + return false; +} + +} // namespace + +bool DeviceCounterEngine::initialize( + const rocprofiler_context_id_t context, + const rocprofiler_agent_id_t gpu_agent, + const uint32_t gpu_device_id, + const MonitorOptions& opts) { + context_ = context; + gpu_agent_ = gpu_agent; + gpu_device_id_ = gpu_device_id; + interval_ns_ = + static_cast(std::max(opts.pm_sampling_interval_us, 1u)) * + 1000u; + max_samples_ = opts.pm_sampling_max_samples; + preset_ = opts.pm_sampling_preset; + pm_sampling_scope_only_ = opts.pm_sampling_scope_only; + deep_arm_mode_ = opts.deep_arm_mode; + sample_row_count_.store(0, std::memory_order_relaxed); + profile_accepted_.store(false, std::memory_order_relaxed); + service_configured_.store(false, std::memory_order_relaxed); + + if (!discoverCounters(gpu_agent_)) { + GFL_LOG_ERROR("[DeviceCounterEngine] no counters discovered"); + return false; + } + + const auto requested = + ResolveAmdDeviceCountingMetrics(opts.pm_sampling_metrics); + if (!createCounterConfig(gpu_agent_, requested)) { + GFL_LOG_ERROR("[DeviceCounterEngine] failed to create counter config"); + return false; + } + + const rocprofiler_buffer_id_t no_buffer{}; + if (!CheckStatus( + rocprofiler_configure_device_counting_service( + context_, no_buffer, gpu_agent_, + &DeviceCounterEngine::configureCallback, this), + "rocprofiler_configure_device_counting_service")) { + shutdown(); + return false; + } + service_configured_.store(true, std::memory_order_release); + + GFL_LOG_DEBUG("[DeviceCounterEngine] initialized device=", gpu_device_id_, + " metrics=", metrics_.size(), + " interval_us=", interval_ns_ / 1000u, + " output_records=", records_.size(), + " scope_gated=", scopeGated() ? "true" : "false"); + return true; +} + +void DeviceCounterEngine::start() { + std::lock_guard lock(sample_mutex_); + session_active_ = true; + + if (!config_emitted_) { + Monitor::EmitPmSamplingConfig( + gpu_device_id_, static_cast(interval_ns_ / 1000u), + max_samples_, preset_, metrics_); + config_emitted_ = true; + } + + if (!profile_accepted_.load(std::memory_order_acquire)) { + // ROCprofiler may invoke the service callback just after + // rocprofiler_start_context returns. session_active_ remains true so + // the first scope can arm once the callback accepts the profile. + GFL_LOG_DEBUG( + "[DeviceCounterEngine] awaiting ROCprofiler counter " + "configuration callback"); + return; + } + if (!scopeGated()) startSamplingLocked(); +} + +void DeviceCounterEngine::stop() { + std::lock_guard lock(sample_mutex_); + stopSamplingLocked(); + session_active_ = false; +} + +void DeviceCounterEngine::drain() { + std::lock_guard lock(sample_mutex_); + sampleLocked(true); +} + +void DeviceCounterEngine::service() { + std::lock_guard lock(sample_mutex_); + sampleLocked(false); +} + +void DeviceCounterEngine::shutdown() { + { + std::lock_guard lock(sample_mutex_); + stopSamplingLocked(); + session_active_ = false; + } + profile_accepted_.store(false, std::memory_order_release); + service_configured_.store(false, std::memory_order_release); + if (config_valid_.exchange(false, std::memory_order_acq_rel)) { + (void) CheckStatus(rocprofiler_destroy_counter_config(config_id_), + "rocprofiler_destroy_counter_config"); + } +} + +void DeviceCounterEngine::onScopeStart(const char*) { + if (!scopeGated()) return; + std::lock_guard lock(sample_mutex_); + if (session_active_) startSamplingLocked(); +} + +void DeviceCounterEngine::onScopeStop(const char*) { + if (!scopeGated()) return; + std::lock_guard lock(sample_mutex_); + stopSamplingLocked(); +} + +bool DeviceCounterEngine::discoverCounters( + const rocprofiler_agent_id_t agent) { + std::vector counter_ids; + const auto callback = + [](rocprofiler_agent_id_t, rocprofiler_counter_id_t* counters, + const size_t count, void* user_data) -> rocprofiler_status_t { + auto* ids = + static_cast*>(user_data); + ids->insert(ids->end(), counters, counters + count); + return ROCPROFILER_STATUS_SUCCESS; + }; + if (!CheckStatus(rocprofiler_iterate_agent_supported_counters( + agent, callback, &counter_ids), + "rocprofiler_iterate_agent_supported_counters")) { + return false; + } + + for (const auto id : counter_ids) { + rocprofiler_counter_info_v0_t info_v0{}; + if (rocprofiler_query_counter_info( + id, ROCPROFILER_COUNTER_INFO_VERSION_0, &info_v0) != + ROCPROFILER_STATUS_SUCCESS) { + continue; + } + + size_t record_count = 1; + rocprofiler_counter_info_v1_t info_v1{}; + if (rocprofiler_query_counter_info( + id, ROCPROFILER_COUNTER_INFO_VERSION_1, &info_v1) == + ROCPROFILER_STATUS_SUCCESS && + info_v1.dimensions_instances_count > 0) { + record_count = + static_cast(info_v1.dimensions_instances_count); + } + + const std::string name = info_v0.name ? info_v0.name : ""; + if (!name.empty()) { + counter_info_[id.handle] = CounterInfo{id, name, record_count}; + } + } + return !counter_info_.empty(); +} + +bool DeviceCounterEngine::createCounterConfig( + const rocprofiler_agent_id_t agent, + const std::vector& requested) { + std::vector candidates; + for (const auto& requested_name : requested) { + const auto it = std::find_if( + counter_info_.begin(), counter_info_.end(), + [&requested_name](const auto& entry) { + return entry.second.name == requested_name; + }); + if (it == counter_info_.end()) { + GFL_LOG_WARN("[DeviceCounterEngine] counter unavailable: ", + requested_name); + continue; + } + candidates.push_back(&it->second); + } + if (candidates.empty()) return false; + + std::vector selected; + selected.reserve(candidates.size()); + for (const auto* info : candidates) selected.push_back(info->id); + + auto status = rocprofiler_create_counter_config( + agent, selected.data(), selected.size(), &config_id_); + std::vector configured; + if (status == ROCPROFILER_STATUS_SUCCESS) { + configured = candidates; + } else { + selected.clear(); + for (const auto* candidate : candidates) { + auto trial = selected; + trial.push_back(candidate->id); + rocprofiler_counter_config_id_t trial_config{}; + if (rocprofiler_create_counter_config( + agent, trial.data(), trial.size(), &trial_config) != + ROCPROFILER_STATUS_SUCCESS) { + GFL_LOG_WARN( + "[DeviceCounterEngine] counter conflicts with active set: ", + candidate->name); + continue; + } + if (!selected.empty()) { + (void) rocprofiler_destroy_counter_config(config_id_); + } + config_id_ = trial_config; + selected = std::move(trial); + configured.push_back(candidate); + } + } + if (configured.empty()) return false; + + size_t record_count = 0; + metrics_.clear(); + for (const auto* info : configured) { + metrics_.push_back(info->name); + record_count += info->record_count; + GFL_LOG_DEBUG("[DeviceCounterEngine] configured counter: ", + info->name, " records=", info->record_count); + } + records_.resize(std::max(record_count, 1)); + config_valid_.store(true, std::memory_order_release); + return true; +} + +void DeviceCounterEngine::startSamplingLocked() { + if (armed_.load(std::memory_order_relaxed) || + !profile_accepted_.load(std::memory_order_acquire)) { + return; + } + next_sample_ns_ = detail::GetTimestampNs(); + attribution_start_ns_ = next_sample_ns_; + last_sample_ns_ = 0; + armed_.store(true, std::memory_order_release); + GFL_LOG_DEBUG("[DeviceCounterEngine] sampling armed"); + Monitor::BeginPmScopeAttribution(next_sample_ns_); + if (deep_arm_mode_ == DeepArmMode::WindowOnly) { + GFL_LOG_INFO("deep window armed: amd.device_counting"); + } +} + +void DeviceCounterEngine::stopSamplingLocked() { + if (!armed_.load(std::memory_order_acquire)) return; + std::optional attributed_ts_ns; + if (scopeGated()) { + // The scope close timestamp is captured immediately before the perf + // stop hook reaches us. Attribute this final pull to the midpoint of + // its collection interval rather than to callback overhead after the + // close; otherwise every short scope's only sample appears unscoped. + const int64_t now_ns = detail::GetTimestampNs(); + const int64_t interval_start_ns = + std::max(attribution_start_ns_, last_sample_ns_); + attributed_ts_ns = + interval_start_ns + (now_ns - interval_start_ns) / 2; + } + sampleLocked(true, attributed_ts_ns); + armed_.store(false, std::memory_order_release); + Monitor::EndPmScopeAttribution(); + attribution_start_ns_ = 0; + last_sample_ns_ = 0; +} + +void DeviceCounterEngine::sampleLocked( + const bool force, const std::optional attributed_ts_ns) { + if (!armed_.load(std::memory_order_acquire) || records_.empty()) return; + + const int64_t before_ns = detail::GetTimestampNs(); + if (!force && before_ns < next_sample_ns_) return; + next_sample_ns_ = + before_ns + static_cast(interval_ns_); + + rocprofiler_user_data_t user_data{}; + user_data.value = sample_index_; + size_t record_count = records_.size(); + auto status = rocprofiler_sample_device_counting_service( + context_, user_data, ROCPROFILER_COUNTER_FLAG_NONE, records_.data(), + &record_count); + if (status == ROCPROFILER_STATUS_ERROR_OUT_OF_RESOURCES && + record_count > records_.size()) { + records_.resize(record_count); + status = rocprofiler_sample_device_counting_service( + context_, user_data, ROCPROFILER_COUNTER_FLAG_NONE, + records_.data(), &record_count); + } + if (status != ROCPROFILER_STATUS_SUCCESS) { + ++sample_failures_; + if (sample_failures_ == 1) { + GFL_LOG_ERROR( + "[DeviceCounterEngine] device counter sample failed: status=", + static_cast(status)); + } + return; + } + records_.resize(record_count); + sample_failures_ = 0; + + std::unordered_map values; + for (size_t i = 0; i < record_count; ++i) { + const auto& record = records_[i]; + rocprofiler_counter_id_t counter_id{}; + const auto query_status = + rocprofiler_query_record_counter_id(record.id, &counter_id); + if (record.agent_id.handle != 0 && + record.agent_id.handle != gpu_agent_.handle) { + continue; + } + if (query_status != ROCPROFILER_STATUS_SUCCESS) { + continue; + } + + std::string counter_name; + if (const auto info = counter_info_.find(counter_id.handle); + info != counter_info_.end()) { + counter_name = info->second.name; + } else { + // ROCprofiler 1.0 on gfx1201 can return the configured derived + // counter under a different high-word generation than the ID + // discovered before HSA startup. Querying the record's own ID is + // authoritative and preserves the native metric name. + rocprofiler_counter_info_v0_t record_info{}; + if (rocprofiler_query_counter_info( + counter_id, ROCPROFILER_COUNTER_INFO_VERSION_0, + &record_info) == ROCPROFILER_STATUS_SUCCESS && + record_info.name) { + counter_name = record_info.name; + counter_info_.emplace( + counter_id.handle, + CounterInfo{counter_id, counter_name, 1}); + } + } + if (counter_name.empty() || !std::isfinite(record.counter_value)) { + continue; + } + values[counter_name] += record.counter_value; + } + + const int64_t after_ns = detail::GetTimestampNs(); + const int64_t sample_ns = attributed_ts_ns.value_or( + before_ns + (after_ns - before_ns) / 2); + std::vector rows; + rows.reserve(metrics_.size()); + for (const auto& metric : metrics_) { + const auto value = values.find(metric); + if (value == values.end()) continue; + PmSampleInput row; + row.sample_index = sample_index_; + row.ts_ns = sample_ns; + row.device_id = gpu_device_id_; + row.metric_name = metric; + row.value = value->second; + rows.push_back(std::move(row)); + } + ++sample_index_; + if (rows.empty()) return; + + last_sample_ns_ = sample_ns; + Monitor::PushPmSamples(rows); + Monitor::PublishScopeRetentionWatermark(sample_ns); + sample_row_count_.fetch_add(rows.size(), std::memory_order_relaxed); +} + +void DeviceCounterEngine::configureCallback( + const rocprofiler_context_id_t context, + const rocprofiler_agent_id_t agent, + const rocprofiler_device_counting_agent_cb_t set_config, + void* user_data) { + auto* engine = static_cast(user_data); + if (!engine || !set_config || agent.handle != engine->gpu_agent_.handle) { + return; + } + const auto status = set_config(context, engine->config_id_); + const bool accepted = status == ROCPROFILER_STATUS_SUCCESS; + engine->profile_accepted_.store(accepted, std::memory_order_release); + if (accepted) { + std::lock_guard lock(engine->sample_mutex_); + if (engine->session_active_ && !engine->scopeGated()) { + engine->startSamplingLocked(); + } + } + if (status != ROCPROFILER_STATUS_SUCCESS) { + GFL_LOG_ERROR( + "[DeviceCounterEngine] set_config failed at context start: status=", + static_cast(status)); + } +} + +} // namespace gpufl::amd diff --git a/include/gpufl/backends/amd/engine/device_counter_engine.hpp b/include/gpufl/backends/amd/engine/device_counter_engine.hpp new file mode 100644 index 0000000..2e2a3eb --- /dev/null +++ b/include/gpufl/backends/amd/engine/device_counter_engine.hpp @@ -0,0 +1,107 @@ +#pragma once + +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include + +#include "gpufl/backends/amd/engine/amd_profiling_engine.hpp" +#include "gpufl/core/monitor.hpp" + +namespace gpufl::amd { + +/// Pulls device-wide hardware counters through ROCprofiler's synchronous +/// device-counting service and emits them through the PM time-series schema. +class DeviceCounterEngine final : public AmdProfilingEngine { + public: + DeviceCounterEngine() = default; + ~DeviceCounterEngine() override { shutdown(); } + + bool initialize(rocprofiler_context_id_t context, + rocprofiler_agent_id_t gpu_agent, + uint32_t gpu_device_id, + const MonitorOptions& opts) override; + void start() override; + void stop() override; + void drain() override; + void service() override; + void shutdown() override; + + bool hasData() const override { + return sample_row_count_.load(std::memory_order_relaxed) > 0; + } + bool isPrepared() const override { + return config_valid_.load(std::memory_order_acquire) && + service_configured_.load(std::memory_order_acquire); + } + bool isArmed() const override { + return armed_.load(std::memory_order_acquire); + } + + void onScopeStart(const char* name) override; + void onScopeStop(const char* name) override; + + private: + struct CounterInfo { + rocprofiler_counter_id_t id{}; + std::string name; + size_t record_count = 1; + }; + + bool discoverCounters(rocprofiler_agent_id_t agent); + bool createCounterConfig(rocprofiler_agent_id_t agent, + const std::vector& requested); + bool scopeGated() const { + return pm_sampling_scope_only_ || + deep_arm_mode_ == DeepArmMode::WindowOnly; + } + void startSamplingLocked(); + void stopSamplingLocked(); + void sampleLocked(bool force, + std::optional attributed_ts_ns = std::nullopt); + + static void configureCallback( + rocprofiler_context_id_t context, + rocprofiler_agent_id_t agent, + rocprofiler_device_counting_agent_cb_t set_config, + void* user_data); + + rocprofiler_context_id_t context_{}; + rocprofiler_agent_id_t gpu_agent_{}; + uint32_t gpu_device_id_ = 0; + rocprofiler_counter_config_id_t config_id_{}; + + std::unordered_map counter_info_; + std::vector metrics_; + std::vector records_; + + uint64_t interval_ns_ = 100'000; + uint32_t max_samples_ = 4096; + std::string preset_ = "overview"; + bool pm_sampling_scope_only_ = true; + DeepArmMode deep_arm_mode_ = DeepArmMode::Always; + + mutable std::mutex sample_mutex_; + bool session_active_ = false; + bool config_emitted_ = false; + int64_t next_sample_ns_ = 0; + int64_t attribution_start_ns_ = 0; + int64_t last_sample_ns_ = 0; + uint32_t sample_index_ = 0; + uint64_t sample_failures_ = 0; + + std::atomic config_valid_{false}; + std::atomic service_configured_{false}; + std::atomic profile_accepted_{false}; + std::atomic armed_{false}; + std::atomic sample_row_count_{0}; +}; + +} // namespace gpufl::amd diff --git a/include/gpufl/backends/amd/rocprofiler_backend.cpp b/include/gpufl/backends/amd/rocprofiler_backend.cpp index b6de107..169c6c4 100644 --- a/include/gpufl/backends/amd/rocprofiler_backend.cpp +++ b/include/gpufl/backends/amd/rocprofiler_backend.cpp @@ -26,6 +26,7 @@ #include #include "gpufl/backends/amd/engine/dispatch_counter_engine.hpp" +#include "gpufl/backends/amd/engine/device_counter_engine.hpp" #include "gpufl/backends/amd/amd_capture_capabilities.hpp" #include "gpufl/core/common.hpp" #include "gpufl/core/debug_logger.hpp" @@ -131,6 +132,7 @@ void RocprofilerBackend::initialize(const MonitorOptions& opts) { capture_capabilities_session_id_.clear(); capability_kernel_rows_baseline_ = 0; capability_memcpy_rows_baseline_ = 0; + capability_pm_sample_rows_baseline_ = 0; capability_dropped_records_baseline_ = 0; capability_queue_dropped_records_baseline_ = 0; capability_buffer_flush_failures_baseline_ = 0; @@ -165,6 +167,10 @@ void RocprofilerBackend::resetToolState() { client_finalize_ = nullptr; tool_registered_.store(false); active_.store(false); + start_requested_.store(false); + next_start_retry_ns_.store(0); + deferred_start_logged_.store(false); + start_failure_logged_.store(false); { std::lock_guard lock(kernel_meta_mutex_); kernel_metadata_.clear(); @@ -320,6 +326,8 @@ int RocprofilerBackend::toolInitialize() { AmdProfilingSupport support; support.dispatch_counting = primary_gpu_agent_.handle != 0 && primary_device_id.has_value(); + support.device_counting = primary_gpu_agent_.handle != 0 && + primary_device_id.has_value(); auto plan = ResolveAmdProfilingPlan(opts_.profiling_engine, support); setResolvedPlan(plan); @@ -342,6 +350,18 @@ int RocprofilerBackend::toolInitialize() { setResolvedPlan( ResolveAmdProfilingPlan(opts_.profiling_engine, support)); } + } else if (plan.selected_path == AmdProfilingPath::DeviceCounting) { + engine_ = std::make_unique(); + if (!engine_->initialize(context_, primary_gpu_agent_, + *primary_device_id, opts_)) { + GFL_LOG_ERROR( + "[ROCProfilerBackend] Device-counter initialization failed; " + "continuing with ROCprofiler trace activity only"); + engine_.reset(); + support.device_counting = false; + setResolvedPlan( + ResolveAmdProfilingPlan(opts_.profiling_engine, support)); + } } tool_registered_.store(true, std::memory_order_release); @@ -360,14 +380,64 @@ void RocprofilerBackend::toolFinalize() { } void RocprofilerBackend::start() { - if (!initialized_.load() || context_.handle == 0 || active_.load()) return; - if (CheckStatus(rocprofiler_start_context(context_), "rocprofiler_start_context")) { - active_.store(true); + if (!initialized_.load(std::memory_order_acquire) || + context_.handle == 0) return; + start_requested_.store(true, std::memory_order_release); + (void) tryStartContext(true); +} + +bool RocprofilerBackend::tryStartContext(const bool force) { + if (!start_requested_.load(std::memory_order_acquire) || + !initialized_.load(std::memory_order_acquire) || + context_.handle == 0 || + active_.load(std::memory_order_acquire)) { + return active_.load(std::memory_order_acquire); + } + + const int64_t now_ns = detail::GetTimestampNs(); + if (!force && + now_ns < next_start_retry_ns_.load(std::memory_order_relaxed)) { + return false; + } + + std::lock_guard lock(start_stop_mutex_); + if (!start_requested_.load(std::memory_order_acquire) || + active_.load(std::memory_order_acquire)) { + return active_.load(std::memory_order_acquire); + } + + const auto status = rocprofiler_start_context(context_); + if (status == ROCPROFILER_STATUS_SUCCESS) { + active_.store(true, std::memory_order_release); + deferred_start_logged_.store(false, std::memory_order_relaxed); + start_failure_logged_.store(false, std::memory_order_relaxed); if (engine_) engine_->start(); + return true; + } + + // gpufl::init commonly runs before the application's first HIP call. + // ROCprofiler cannot start until that call loads HSA, so retain the + // configured context and retry from the collector or first scope. + next_start_retry_ns_.store(now_ns + 10'000'000, + std::memory_order_relaxed); + if (status == ROCPROFILER_STATUS_ERROR_HSA_NOT_LOADED) { + if (!deferred_start_logged_.exchange(true, + std::memory_order_relaxed)) { + GFL_LOG_DEBUG( + "[ROCProfilerBackend] deferring context start until HSA is " + "loaded"); + } + } else if (!start_failure_logged_.exchange( + true, std::memory_order_relaxed)) { + GFL_LOG_ERROR("[ROCProfilerBackend] rocprofiler_start_context failed: ", + StatusToString(status)); } + return false; } void RocprofilerBackend::stop() { + start_requested_.store(false, std::memory_order_release); + std::lock_guard lock(start_stop_mutex_); if (!active_.exchange(false) || context_.handle == 0) return; if (engine_) engine_->stop(); (void) CheckStatus(rocprofiler_stop_context(context_), @@ -387,12 +457,11 @@ void RocprofilerBackend::DrainProfilingData() { } void RocprofilerBackend::ServiceDeepWindow() { - if (!initialized_.load(std::memory_order_acquire) || - !active_.load(std::memory_order_acquire) || - !DeepWindow::HasPendingWork()) { - return; - } - DeepWindow::ServicePending(); + if (!initialized_.load(std::memory_order_acquire)) return; + if (!active_.load(std::memory_order_acquire) && + !tryStartContext(false)) return; + if (DeepWindow::HasPendingWork()) DeepWindow::ServicePending(); + if (engine_) engine_->service(); } bool RocprofilerBackend::DeepEnginesPrepared() const { @@ -410,6 +479,9 @@ std::vector RocprofilerBackend::OnDeepWindowStop( } void RocprofilerBackend::OnPerfScopeStart(const char* name) { + if (!active_.load(std::memory_order_acquire)) { + (void) tryStartContext(true); + } if (opts_.deep_arm_mode == DeepArmMode::WindowOnly) return; OnDeepWindowPerfStart(name); } @@ -420,6 +492,9 @@ void RocprofilerBackend::OnPerfScopeStop(const char* name) { } void RocprofilerBackend::OnDeepWindowPerfStart(const char* name) { + if (!active_.load(std::memory_order_acquire)) { + (void) tryStartContext(true); + } if (engine_) engine_->onScopeStart(name); } @@ -453,6 +528,7 @@ void RocprofilerBackend::emitCapabilities() { trace_records_unattributed_.load(std::memory_order_relaxed); const uint64_t scope_correlation_failures = scope_correlation_failures_.load(std::memory_order_relaxed); + const uint64_t pm_sample_rows = Monitor::PmSampleRowsSeen(); AmdCaptureCapabilityInput input; input.session_id = segment->session_id; @@ -466,6 +542,8 @@ void RocprofilerBackend::emitCapabilities() { delta(memcpy_rows, capability_memcpy_rows_baseline_); input.profiling_sample_rows = engine_ && engine_->hasData() ? 1 : 0; + input.pm_sample_rows = + delta(pm_sample_rows, capability_pm_sample_rows_baseline_); input.dropped_trace_records = delta(dropped_records, capability_dropped_records_baseline_); input.dropped_client_records = delta( @@ -483,6 +561,7 @@ void RocprofilerBackend::emitCapabilities() { capability_kernel_rows_baseline_ = kernel_rows; capability_memcpy_rows_baseline_ = memcpy_rows; + capability_pm_sample_rows_baseline_ = pm_sample_rows; capability_dropped_records_baseline_ = dropped_records; capability_queue_dropped_records_baseline_ = queue_dropped_records; capability_buffer_flush_failures_baseline_ = buffer_flush_failures; @@ -524,7 +603,9 @@ bool RocprofilerBackend::flushBuffers() { } void RocprofilerBackend::OnScopeStart(const char* name) { - if (!active_.load() || context_.handle == 0 || name == nullptr) return; + if (context_.handle == 0 || name == nullptr) return; + if (!active_.load(std::memory_order_acquire) && + !tryStartContext(true)) return; rocprofiler_thread_id_t tid{}; if (!CheckStatus(rocprofiler_get_thread_id(&tid), diff --git a/include/gpufl/backends/amd/rocprofiler_backend.hpp b/include/gpufl/backends/amd/rocprofiler_backend.hpp index 82d2be7..9c9db66 100644 --- a/include/gpufl/backends/amd/rocprofiler_backend.hpp +++ b/include/gpufl/backends/amd/rocprofiler_backend.hpp @@ -75,6 +75,7 @@ class RocprofilerBackend final : public IMonitorBackend { bool configureRocprofiler(const MonitorOptions& opts, std::string* reason); void resetToolState(); bool registerTool(std::string* reason); + bool tryStartContext(bool force); AmdResolvedProfilingPlan resolvedPlan() const; void setResolvedPlan(AmdResolvedProfilingPlan plan); @@ -163,6 +164,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_pm_sample_rows_baseline_ = 0; mutable uint64_t capability_dropped_records_baseline_ = 0; mutable uint64_t capability_queue_dropped_records_baseline_ = 0; mutable uint64_t capability_buffer_flush_failures_baseline_ = 0; @@ -171,6 +173,11 @@ class RocprofilerBackend final : public IMonitorBackend { std::atomic initialized_{false}; std::atomic active_{false}; + std::atomic start_requested_{false}; + std::atomic next_start_retry_ns_{0}; + std::atomic deferred_start_logged_{false}; + std::atomic start_failure_logged_{false}; + std::mutex start_stop_mutex_; std::atomic tool_registered_{false}; }; diff --git a/include/gpufl/report/text_report.cpp b/include/gpufl/report/text_report.cpp index 5601c8c..47b0dd0 100644 --- a/include/gpufl/report/text_report.cpp +++ b/include/gpufl/report/text_report.cpp @@ -77,6 +77,18 @@ std::string shortenKernelName(const std::string& name) { std::string s = name; auto at = s.find('@'); if (at != std::string::npos) s = s.substr(0, at); + + // GCC/ROCm demangling prefixes internal kernels with a parenthesized + // namespace, for example `(anonymous namespace)::kernel(float*)`. The + // argument-list scan below would otherwise mistake that first `(` for + // the function parameters and shorten the display name to an empty + // string. + constexpr const char* kAnonymousNamespace = "(anonymous namespace)::"; + for (auto pos = s.find(kAnonymousNamespace); pos != std::string::npos; + pos = s.find(kAnonymousNamespace)) { + s.erase(pos, std::char_traits::length(kAnonymousNamespace)); + } + for (const auto* prefix : {"void ", "int ", "float ", "double ", "__global__ "}) { if (s.rfind(prefix, 0) == 0) { s = s.substr(std::string(prefix).size()); @@ -353,6 +365,7 @@ void TextReport::parseJobStart(const JsonValue& rec, SessionInfo& info) { if (!rec.contains(field) || !rec[field].is_array() || rec[field].empty()) continue; const auto& dev = rec[field][0]; info.gpu_name = dev.value("name", ""); + info.gpu_vendor = dev.value("vendor", ""); info.compute_major = dev.contains("compute_capability_major") ? dev.value("compute_capability_major", 0) : dev.value("major", 0); info.compute_minor = dev.contains("compute_capability_minor") @@ -365,8 +378,23 @@ void TextReport::parseJobStart(const JsonValue& rec, SessionInfo& info) { } if (info.gpu_name.empty() && rec.contains("devices") && - rec["devices"].is_array() && !rec["devices"].empty()) + rec["devices"].is_array() && !rec["devices"].empty()) { info.gpu_name = rec["devices"][0].value("name", ""); + info.gpu_vendor = rec["devices"][0].value("vendor", ""); + } +} + +bool TextReport::isAmdSession() const { + if (selected_engine_.rfind("amd.", 0) == 0) return true; + + std::string identity = info_.gpu_vendor + " " + info_.gpu_name; + std::transform(identity.begin(), identity.end(), identity.begin(), + [](const unsigned char ch) { + return static_cast(std::tolower(ch)); + }); + return identity.find("amd") != std::string::npos || + identity.find("radeon") != std::string::npos || + identity.find("advanced micro devices") != std::string::npos; } void TextReport::parseDeviceLog(const std::vector& records, @@ -639,11 +667,15 @@ void TextReport::writeSessionSummary(std::ostringstream& out) const { out << " Duration: " << fmtDuration((info_.end_ns - info_.start_ns) / 1e6) << "\n"; if (!info_.gpu_name.empty()) { + const bool amd = isAmdSession(); out << " GPU Device: " << info_.gpu_name << "\n"; if (info_.compute_major > 0) out << " Compute: " << info_.compute_major << "." << info_.compute_minor << "\n"; - if (info_.sm_count > 0) - out << " SMs: " << info_.sm_count << "\n"; + if (info_.sm_count > 0) { + out << " " << std::left << std::setw(20) + << (amd ? "Compute Units:" : "SMs:") << info_.sm_count + << std::right << "\n"; + } if (info_.shared_mem_per_block > 0) out << " Shared Mem/Block: " << fmtBytes(info_.shared_mem_per_block) << "\n"; if (info_.regs_per_block > 0) @@ -792,6 +824,7 @@ void TextReport::writeKernelDetails(std::ostringstream& out) const { for (const auto& k : kernels_) totalByName[k.name] += k.duration_ms; auto ranked = sortedTopN(totalByName, top_n_, [](double v) { return v; }); + const bool amd = isAmdSession(); for (const auto& [name, _] : ranked) { // Find representative kernel with detail data auto rep = std::find_if(kernels_.begin(), kernels_.end(), @@ -813,8 +846,10 @@ void TextReport::writeKernelDetails(std::ostringstream& out) const { << std::fixed << std::setprecision(1) << (val * 100) << "%\n"; }; writeOcc("Reg Occupancy", rep->reg_occupancy); - writeOcc("SMem Occupancy", rep->smem_occupancy); - writeOcc("Warp Occupancy", rep->warp_occupancy); + writeOcc(amd ? "LDS Occupancy" : "SMem Occupancy", + rep->smem_occupancy); + writeOcc(amd ? "Wave Occupancy" : "Warp Occupancy", + rep->warp_occupancy); writeOcc("Block Occupancy", rep->block_occupancy); if (!rep->limiting_resource.empty()) @@ -836,7 +871,9 @@ void TextReport::writeKernelDetails(std::ostringstream& out) const { double waves = static_cast(gridTotal) / (static_cast(rep->max_active_blocks) * info_.sm_count); - out << " Waves/SM: " << std::fixed + out << (amd ? " Waves/CU: " + : " Waves/SM: ") + << std::fixed << std::setprecision(2) << waves << "\n"; } } @@ -1084,7 +1121,7 @@ void TextReport::writeScopeSummary(std::ostringstream& out) const { if (!scopeGpu.empty()) { auto ranked = sortedTopN(scopeGpu, 0, [](const AggStats& s) { return s.total; }); - out << "\n GPU Time by Scope (kernel execution only - SM time from CUPTI):\n"; + out << "\n GPU Time by Scope (kernel execution only):\n"; out << " " << std::left << std::setw(30) << "Scope" << std::right << std::setw(8) << "Kernels" << std::setw(14) << "GPU Time" << std::setw(12) << "Avg" << "\n"; @@ -1102,7 +1139,7 @@ void TextReport::writeScopeSummary(std::ostringstream& out) const { if (!scopeGpu.empty()) { out << "\n Note: Scope Timing is what your CPU thread saw " "(includes JIT, sync, launch overhead).\n"; - out << " GPU Time is pure kernel execution on the SM. " + out << " GPU Time is pure kernel execution on the GPU. " "Use GPU Time to compare kernel perf.\n"; } } @@ -1262,17 +1299,31 @@ static std::string makeBar(double pct, int maxWidth = 20) { using FuncProfile = gpufl::report::FuncProfile; void TextReport::writeProfileAnalysis(std::ostringstream& out) const { - out << "\n" << SEP << "\n Profile / SASS Analysis\n" << SEP << "\n"; - if (profile_samples_.empty()) { out << " (No profile sample data)\n"; return; } + // `isa_inst_present` rows describe disassembly/source-map availability; + // they are metadata rather than performance samples. Do not create an + // otherwise empty instruction-analysis section for trace or PM-only runs. + const bool hasAnalysisSample = std::any_of( + profile_samples_.begin(), profile_samples_.end(), + [](const ProfileSampleRecord& ps) { + return ps.metric_name != "isa_inst_present"; + }); + if (!hasAnalysisSample) return; + + out << "\n" << SEP << "\n Profile / Instruction Analysis\n" << SEP + << "\n"; const bool hasNonZeroSample = std::any_of( profile_samples_.begin(), profile_samples_.end(), - [](const ProfileSampleRecord& ps) { return ps.metric_value > 0; }); + [](const ProfileSampleRecord& ps) { + return ps.metric_name != "isa_inst_present" && + ps.metric_value > 0; + }); if (!hasNonZeroSample) { out << " (Profile sample rows were present, but every metric value was 0.)\n"; if (sass_active_) { - out << " SASS instrumentation was configured, but CUPTI returned no " - "non-zero SASS counter values for this session.\n"; + out << " Instruction instrumentation was configured, but the " + "profiler returned no non-zero counter values for this " + "session.\n"; } return; } @@ -1322,6 +1373,7 @@ void TextReport::writeProfileAnalysis(std::ostringstream& out) const { // ── Collect per-function data ─────────────────────────────────────────── std::map byFunc; for (const auto& ps : profile_samples_) { + if (ps.metric_name == "isa_inst_present") continue; std::string fn = ps.function_name.empty() ? "(unknown)" : ps.function_name; auto& fp = byFunc[fn]; @@ -1425,6 +1477,7 @@ void TextReport::writeProfileAnalysis(std::ostringstream& out) const { std::map otherMetrics; for (const auto& ps : profile_samples_) { if (ps.metric_name.empty()) continue; + if (ps.metric_name == "isa_inst_present") continue; if (ps.stall_reason > 1 && ps.sample_kind == 0) continue; // stall data already shown if (ps.metric_name == "smsp__sass_inst_executed") continue; if (ps.metric_name == "smsp__sass_thread_inst_executed") continue; @@ -1437,7 +1490,7 @@ void TextReport::writeProfileAnalysis(std::ostringstream& out) const { if (!otherMetrics.empty()) { auto metricRanked = sortedTopN(otherMetrics, 0, [](uint64_t v) { return static_cast(v); }); - out << "\n Other SASS Metrics:\n"; + out << "\n Other Profile Metrics:\n"; out << " " << std::left << std::setw(50) << "Metric" << std::right << std::setw(16) << "Total" << "\n"; out << " " << std::string(66, '-') << "\n"; diff --git a/include/gpufl/report/text_report.hpp b/include/gpufl/report/text_report.hpp index c5caf06..aa5b1c2 100644 --- a/include/gpufl/report/text_report.hpp +++ b/include/gpufl/report/text_report.hpp @@ -139,6 +139,7 @@ class TextReport { int64_t start_ns = 0; int64_t end_ns = 0; std::string gpu_name; + std::string gpu_vendor; int compute_major = 0; int compute_minor = 0; int sm_count = 0; @@ -199,6 +200,7 @@ class TextReport { void parseSystemLog(const std::vector& records); void parseCaptureCapabilities(const std::vector& records); void mergeKernelDetails(std::unordered_map& details); + bool isAmdSession() const; static void parseJobStart(const JsonValue& record, SessionInfo& info); diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index ae77592..2216a7e 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -38,6 +38,7 @@ set(GPUFL_TEST_SOURCES core/test_session_bootstrap.cpp core/test_source_capture_policy.cpp core/test_startup_configuration.cpp + core/test_text_report.cpp upload/test_upload_logs.cpp # Launcher CLI parser test - portable (no CUDA / no POSIX). # The portable CLI parser sources are compiled directly into the test diff --git a/tests/backends/amd/test_amd_profiling_policy.cpp b/tests/backends/amd/test_amd_profiling_policy.cpp index 11c8ca8..744d67f 100644 --- a/tests/backends/amd/test_amd_profiling_policy.cpp +++ b/tests/backends/amd/test_amd_profiling_policy.cpp @@ -88,6 +88,17 @@ TEST(AmdProfilingPolicy, RequestIntentNeverInventsAmdNativeNames) { "amd.device_counting"); } +TEST(AmdProfilingPolicy, DeviceCountingUsesPortableDefaultMetric) { + EXPECT_EQ(gpufl::amd::ResolveAmdDeviceCountingMetrics({}), + std::vector{"GPUBusy"}); +} + +TEST(AmdProfilingPolicy, DeviceCountingPreservesUniqueExplicitMetrics) { + EXPECT_EQ(gpufl::amd::ResolveAmdDeviceCountingMetrics( + {"SQ_WAVES", "", "GPUBusy", "SQ_WAVES"}), + (std::vector{"SQ_WAVES", "GPUBusy"})); +} + TEST(AmdProfilingPolicy, TraceSelectsBufferTracingService) { const auto plan = gpufl::amd::ResolveAmdProfilingPlan( gpufl::ProfilingEngine::Trace, {}); @@ -211,6 +222,34 @@ TEST(AmdCaptureCapabilities, DispatchSamplesAndDroppedTraceAreVisible) { std::string::npos); } +TEST(AmdCaptureCapabilities, DeviceCountingUsesPmRowsForDataStatus) { + gpufl::amd::AmdProfilingSupport support; + support.device_counting = true; + + gpufl::amd::AmdCaptureCapabilityInput input; + input.session_id = "amd-session"; + input.plan = gpufl::amd::ResolveAmdProfilingPlan( + gpufl::ProfilingEngine::PmSampling, support); + input.trace_configured = true; + input.profiling_sample_rows = 0; + input.pm_sample_rows = 3; + + const auto event = gpufl::amd::BuildAmdCaptureCapabilitiesEvent(input); + const auto* counters = FindCapability(event, "device_counting"); + ASSERT_NE(counters, nullptr); + EXPECT_TRUE(counters->requested); + EXPECT_EQ(counters->status, "collected"); + EXPECT_EQ(counters->mode, "rocprofiler_device_counting_service"); + + input.profiling_sample_rows = 3; + input.pm_sample_rows = 0; + const auto no_pm_event = + gpufl::amd::BuildAmdCaptureCapabilitiesEvent(input); + counters = FindCapability(no_pm_event, "device_counting"); + ASSERT_NE(counters, nullptr); + EXPECT_EQ(counters->status, "enabled_no_data"); +} + TEST(AmdCaptureCapabilities, LifecycleDeliveryAndCorrelationFailuresAreVisible) { gpufl::amd::AmdCaptureCapabilityInput input; input.session_id = "amd-session"; diff --git a/tests/core/test_text_report.cpp b/tests/core/test_text_report.cpp new file mode 100644 index 0000000..2e50353 --- /dev/null +++ b/tests/core/test_text_report.cpp @@ -0,0 +1,96 @@ +#include + +#include +#include +#include +#include + +#include "gpufl/report/text_report.hpp" + +namespace fs = std::filesystem; + +namespace { + +class TextReportTest : public ::testing::Test { + protected: + void SetUp() override { + const auto* info = + ::testing::UnitTest::GetInstance()->current_test_info(); + log_dir_ = fs::temp_directory_path() / + (std::string("gpufl_text_report_") + info->name()); + fs::remove_all(log_dir_); + fs::create_directories(log_dir_); + } + + void TearDown() override { + std::error_code ec; + fs::remove_all(log_dir_, ec); + } + + void WriteLog(const std::string& channel, + const std::vector& records) const { + std::ofstream out(log_dir_ / ("fixture." + channel + ".log"), + std::ios::binary | std::ios::trunc); + ASSERT_TRUE(out.good()); + for (const auto& record : records) out << record << '\n'; + } + + std::string Generate() const { + gpufl::report::TextReport::Options options; + options.log_dir = log_dir_.string(); + options.log_prefix = "fixture"; + return gpufl::report::TextReport(options).generate(); + } + + fs::path log_dir_; +}; + +TEST_F(TextReportTest, AmdPmReportUsesKernelNameAndBackendAwareLabels) { + WriteLog("device", { + R"({"type":"job_start","session_id":"s1","app":"amd_report","ts_ns":1000,"gpu_static_devices":[{"name":"AMD Radeon Test","vendor":"AMD","multi_processor_count":32}]})", + R"({"type":"capture_capabilities","session_id":"s1","requested_engine":"pm_sampling","selected_engine":"amd.device_counting","capabilities":[]})", + R"({"type":"dictionary_update","session_id":"s1","kernel_dict":{"1":"(anonymous namespace)::sampleRowsWorkload(float*, int, int) [clone .kd]"}})", + R"({"type":"kernel_event_batch","session_id":"s1","base_time_ns":1000,"columns":["dt_ns","duration_ns","kernel_id","stream_id","corr_id","num_regs","dyn_shared","has_details"],"rows":[[100,1000,1,0,7,136,0,1]]})", + R"json({"type":"kernel_detail","session_id":"s1","corr_id":7,"grid":"(4096,1,1)","block":"(256,1,1)","occupancy":1.0,"reg_occupancy":1.0,"smem_occupancy":1.0,"warp_occupancy":1.0,"block_occupancy":1.0,"limiting_resource":"waves","max_active_blocks":4,"user_scope":"pm_rows_phase_a"})json", + R"({"type":"shutdown","session_id":"s1","ts_ns":3000})", + }); + WriteLog("scope", { + R"({"type":"dictionary_update","session_id":"s1","function_dict":{"1":"sampleRowsWorkload"},"metric_dict":{"1":"isa_inst_present"}})", + R"({"type":"profile_sample_batch","session_id":"s1","columns":["function_id","metric_id","metric_value","stall_reason","sample_kind"],"rows":[[1,1,49,0,1]]})", + }); + + const std::string report = Generate(); + EXPECT_NE(report.find("sampleRowsWorkload"), std::string::npos); + EXPECT_NE(report.find("Compute Units:"), std::string::npos); + EXPECT_NE(report.find("LDS Occupancy:"), std::string::npos); + EXPECT_NE(report.find("Wave Occupancy:"), std::string::npos); + EXPECT_NE(report.find("Waves/CU:"), std::string::npos); + EXPECT_NE(report.find("GPU Time by Scope (kernel execution only):"), + std::string::npos); + EXPECT_EQ(report.find("SM time from CUPTI"), std::string::npos); + EXPECT_EQ(report.find("Profile / SASS Analysis"), std::string::npos); + EXPECT_EQ(report.find("Profile / Instruction Analysis"), + std::string::npos); + EXPECT_EQ(report.find("isa_inst_present"), std::string::npos); +} + +TEST_F(TextReportTest, MeaningfulProfileRowsUseBackendNeutralSectionNames) { + WriteLog("device", { + R"({"type":"job_start","session_id":"s1","app":"amd_dispatch","ts_ns":1000,"gpu_static_devices":[{"name":"AMD Radeon Test","vendor":"AMD","multi_processor_count":32}]})", + R"({"type":"capture_capabilities","session_id":"s1","requested_engine":"sass_metrics","selected_engine":"amd.dispatch_counting","capabilities":[]})", + R"({"type":"shutdown","session_id":"s1","ts_ns":3000})", + }); + WriteLog("scope", { + R"({"type":"dictionary_update","session_id":"s1","function_dict":{"1":"dispatchKernel"},"metric_dict":{"1":"SQ_WAVES"}})", + R"({"type":"profile_sample_batch","session_id":"s1","columns":["function_id","metric_id","metric_value","stall_reason","sample_kind"],"rows":[[1,1,9,0,1]]})", + }); + + const std::string report = Generate(); + EXPECT_NE(report.find("Profile / Instruction Analysis"), + std::string::npos); + EXPECT_NE(report.find("Other Profile Metrics:"), std::string::npos); + EXPECT_NE(report.find("SQ_WAVES"), std::string::npos); + EXPECT_EQ(report.find("Other SASS Metrics:"), std::string::npos); +} + +} // namespace