From 57ac18bc8313f58409a80c650cbb851e776df6d2 Mon Sep 17 00:00:00 2001 From: Marcos Slomp Date: Wed, 24 Jun 2026 13:32:29 -0700 Subject: [PATCH] API exploration for kernel metrics --- public/tracy/TracyCUDA.hpp | 319 +++++++++++++++++++++++++++++++++++++ 1 file changed, 319 insertions(+) diff --git a/public/tracy/TracyCUDA.hpp b/public/tracy/TracyCUDA.hpp index c5134807..b5f3c19c 100644 --- a/public/tracy/TracyCUDA.hpp +++ b/public/tracy/TracyCUDA.hpp @@ -10,6 +10,10 @@ #define TracyCUDAStartProfiling(ctx) #define TracyCUDAStopProfiling(ctx) +#define TracyCUDAEnableKernelMetrics(ctx) +#define TracyCUDADisableKernelMetrics(ctx) +#define TracyCUDADumpKernelMetrics(ctx) + #define TracyCUDACollect(ctx) namespace tracy{ @@ -57,6 +61,24 @@ using CUDACtx = std::nullptr_t; #define TRACY_CUDA_ENABLE_CUDA_CALL_STATS (0) #endif//TRACY_CUDA_ENABLE_CUDA_CALL_STATS +// Opt-in CUPTI Range Profiler support: collects hardware metrics for kernel +// dispatches / graph launches and surfaces them as Tracy plots. OFF by default +// because it pulls in an extra dependency (the NVPW host library) and may +// require elevated GPU performance-counter permissions at runtime. +#ifndef TRACY_CUDA_ENABLE_KERNEL_METRICS +#define TRACY_CUDA_ENABLE_KERNEL_METRICS (0) +#endif//TRACY_CUDA_ENABLE_KERNEL_METRICS + +#if TRACY_CUDA_ENABLE_KERNEL_METRICS +#if CUDA_VERSION < 12060 +#error "kernel metrics require CUDA v12.6 (or later)" +#endif +#include +#include +#include +#include +#endif//TRACY_CUDA_ENABLE_KERNEL_METRICS + namespace { // TODO(marcos): wrap these in structs for better type safety @@ -568,6 +590,22 @@ namespace tracy printStats(); } + void EnableKernelMetrics() + { + ZoneScoped; + #if TRACY_CUDA_ENABLE_KERNEL_METRICS + CUPTI::BeginKernelMetrics(this); + #endif//TRACY_CUDA_ENABLE_KERNEL_METRICS + } + + void DisableKernelMetrics() + { + ZoneScoped; + #if TRACY_CUDA_ENABLE_KERNEL_METRICS + CUPTI::EndKernelMetrics(); + #endif//TRACY_CUDA_ENABLE_KERNEL_METRICS + } + void Name(const char *name, uint16_t len) { auto ptr = (char*)tracyMalloc(len); @@ -1263,6 +1301,262 @@ namespace tracy //CUPTI_ACTIVITY_KIND_DRIVER, }; + #if TRACY_CUDA_ENABLE_KERNEL_METRICS + // Curated set of PerfWorks metrics collected per range (dispatch / graph + // launch). Kept deliberately small so the whole set fits in a SINGLE + // collection pass (no kernel replay): all are high-level "% of peak" + // rollups, so each maps to only a few raw counters. BeginKernelMetrics + // asserts the resulting config needs exactly one pass; trim this list if + // a given architecture's counter budget is exceeded. Note that some + // metric names can vary across GPU architectures. + static constexpr const char* kKernelMetricNames[] = { + // Compute efficiency + "sm__throughput.avg.pct_of_peak_sustained_elapsed", // overall SM/compute throughput + "smsp__inst_executed.avg.per_cycle_active", // instructions issued per active cycle (IPC) + // Memory utilization + "gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed", // overall memory-subsystem throughput + "dram__throughput.avg.pct_of_peak_sustained_elapsed", // device-memory (DRAM) bandwidth + "lts__t_sector_hit_rate.pct", // L2 cache hit rate + // Occupancy / latency hiding + "sm__warps_active.avg.pct_of_peak_sustained_active", // achieved occupancy + // Stalls (issue-slot utilization: low value => warps stalled) + "smsp__issue_active.avg.pct_of_peak_sustained_active", + }; + static constexpr size_t kKernelMetricCount = sizeof(kKernelMetricNames) / sizeof(kKernelMetricNames[0]); + + // Upper bound on ranges (dispatches / graph launches) buffered in one + // counter-data image before it must be decoded and reset. Tunable; sizes + // the counter-data allocation and caps ranges-per-pass at SetConfig time. + static constexpr size_t kMaxRangesPerPass = 64; + + // Range Profiler bring-up: one-time process-wide profiler init plus + // creation of a Range Profiler object bound to the current CUDA context. + // Metric configuration (chip name + curated single-pass metric list -> + // config image), cuptiRangeProfilerSetConfig() in user-range mode, and + // cuptiRangeProfilerStart() are wired up in a later step. + static void BeginKernelMetrics(CUDACtx* profilerHost) { + ZoneScoped; + UNREFERENCED(profilerHost); + auto& rangeProfiler = PersistentState::Get().rangeProfiler; + if (rangeProfiler != nullptr) { + return; // already initialized + } + + // Process-wide CUPTI profiler initialization (idempotent). + CUpti_Profiler_Initialize_Params initParams = { CUpti_Profiler_Initialize_Params_STRUCT_SIZE }; + CUPTI_API_CALL(cuptiProfilerInitialize(&initParams)); + + // Bind a Range Profiler object to the calling thread's current + // context (a null ctx tells CUPTI to use the current context). + CUcontext cuCtx = nullptr; + DRIVER_API_CALL(cuCtxGetCurrent(&cuCtx)); + CUpti_RangeProfiler_Enable_Params createParams = { CUpti_RangeProfiler_Enable_Params_STRUCT_SIZE }; + createParams.ctx = cuCtx; + CUPTI_API_CALL(cuptiRangeProfilerEnable(&createParams)); + rangeProfiler = createParams.pRangeProfilerObject; + + // Query the device chip name: the host profiler is chip-specific and + // needs it to resolve metric names into hardware counters. + CUdevice cuDevice = 0; + DRIVER_API_CALL(cuCtxGetDevice(&cuDevice)); + CUpti_Device_GetChipName_Params chipNameParams = { CUpti_Device_GetChipName_Params_STRUCT_SIZE }; + chipNameParams.deviceIndex = (size_t)cuDevice; + CUPTI_API_CALL(cuptiDeviceGetChipName(&chipNameParams)); + auto& chipName = PersistentState::Get().chipName; + chipName = chipNameParams.pChipName; + + // Snapshot which counters are actually available on this context so + // the host can reject unsupported metrics up front. The first call + // reports the image size; the second fills the allocated buffer. + auto& counterAvailabilityImage = PersistentState::Get().counterAvailabilityImage; + CUpti_Profiler_GetCounterAvailability_Params availabilityParams = { CUpti_Profiler_GetCounterAvailability_Params_STRUCT_SIZE }; + availabilityParams.ctx = cuCtx; + CUPTI_API_CALL(cuptiProfilerGetCounterAvailability(&availabilityParams)); + counterAvailabilityImage.resize(availabilityParams.counterAvailabilityImageSize); + availabilityParams.pCounterAvailabilityImage = counterAvailabilityImage.data(); + CUPTI_API_CALL(cuptiProfilerGetCounterAvailability(&availabilityParams)); + + // Create the host profiler object: it builds metric config images + // and later evaluates counter data into metric values. + CUpti_Profiler_Host_Initialize_Params hostInitParams = { CUpti_Profiler_Host_Initialize_Params_STRUCT_SIZE }; + hostInitParams.profilerType = CUPTI_PROFILER_TYPE_RANGE_PROFILER; + hostInitParams.pChipName = chipName.c_str(); + hostInitParams.pCounterAvailabilityImage = counterAvailabilityImage.data(); + CUPTI_API_CALL(cuptiProfilerHostInitialize(&hostInitParams)); + auto* hostObject = hostInitParams.pHostObject; + PersistentState::Get().profilerHostObject = hostObject; + + // Add the curated metric list to the host config and bake a config + // image describing the hardware counters to collect. + CUpti_Profiler_Host_ConfigAddMetrics_Params addMetricsParams = { CUpti_Profiler_Host_ConfigAddMetrics_Params_STRUCT_SIZE }; + addMetricsParams.pHostObject = hostObject; + addMetricsParams.ppMetricNames = const_cast(kKernelMetricNames); + addMetricsParams.numMetrics = kKernelMetricCount; + CUPTI_API_CALL(cuptiProfilerHostConfigAddMetrics(&addMetricsParams)); + + auto& configImage = PersistentState::Get().configImage; + CUpti_Profiler_Host_GetConfigImageSize_Params configSizeParams = { CUpti_Profiler_Host_GetConfigImageSize_Params_STRUCT_SIZE }; + configSizeParams.pHostObject = hostObject; + CUPTI_API_CALL(cuptiProfilerHostGetConfigImageSize(&configSizeParams)); + configImage.resize(configSizeParams.configImageSize); + CUpti_Profiler_Host_GetConfigImage_Params configImageParams = { CUpti_Profiler_Host_GetConfigImage_Params_STRUCT_SIZE }; + configImageParams.pHostObject = hostObject; + configImageParams.pConfigImage = configImage.data(); + configImageParams.configImageSize = configImage.size(); + CUPTI_API_CALL(cuptiProfilerHostGetConfigImage(&configImageParams)); + + // Enforce the single-pass (no-replay) design constraint up front: if + // the selected metrics don't fit one pass, trim kKernelMetricNames. + CUpti_Profiler_Host_GetNumOfPasses_Params numPassesParams = { CUpti_Profiler_Host_GetNumOfPasses_Params_STRUCT_SIZE }; + numPassesParams.pConfigImage = configImage.data(); + numPassesParams.configImageSize = configImage.size(); + CUPTI_API_CALL(cuptiProfilerHostGetNumOfPasses(&numPassesParams)); + if (numPassesParams.numOfPasses != 1) { + fprintf(stderr, "ERROR:\tTracyCUDA kernel metrics need a single-pass config, " + "but the selected metrics require %llu passes; trim kKernelMetricNames.\n", + (unsigned long long)numPassesParams.numOfPasses); + assert(numPassesParams.numOfPasses == 1); + } + + // Allocate and initialize the counter-data image the Range Profiler + // fills during collection (sized for the metric set and range batch). + auto& counterDataImage = PersistentState::Get().counterDataImage; + CUpti_RangeProfiler_GetCounterDataSize_Params counterDataSizeParams = { CUpti_RangeProfiler_GetCounterDataSize_Params_STRUCT_SIZE }; + counterDataSizeParams.pRangeProfilerObject = rangeProfiler; + counterDataSizeParams.pMetricNames = const_cast(kKernelMetricNames); + counterDataSizeParams.numMetrics = kKernelMetricCount; + counterDataSizeParams.maxNumOfRanges = kMaxRangesPerPass; + counterDataSizeParams.maxNumRangeTreeNodes = kMaxRangesPerPass; + CUPTI_API_CALL(cuptiRangeProfilerGetCounterDataSize(&counterDataSizeParams)); + counterDataImage.resize(counterDataSizeParams.counterDataSize); + CUpti_RangeProfiler_CounterDataImage_Initialize_Params counterDataInitParams = { CUpti_RangeProfiler_CounterDataImage_Initialize_Params_STRUCT_SIZE }; + counterDataInitParams.pRangeProfilerObject = rangeProfiler; + counterDataInitParams.counterDataSize = counterDataImage.size(); + counterDataInitParams.pCounterData = counterDataImage.data(); + CUPTI_API_CALL(cuptiRangeProfilerCounterDataImageInitialize(&counterDataInitParams)); + + // Bind config + counter-data to the Range Profiler in user-range mode. + // A single pass means the user-replay loop runs exactly once, so there + // is no kernel re-execution and no forced serialization of overlap. + CUpti_RangeProfiler_SetConfig_Params setConfigParams = { CUpti_RangeProfiler_SetConfig_Params_STRUCT_SIZE }; + setConfigParams.pRangeProfilerObject = rangeProfiler; + setConfigParams.pConfig = configImage.data(); + setConfigParams.configSize = configImage.size(); + setConfigParams.pCounterDataImage = counterDataImage.data(); + setConfigParams.counterDataImageSize = counterDataImage.size(); + setConfigParams.range = CUPTI_UserRange; + setConfigParams.replayMode = CUPTI_UserReplay; + setConfigParams.maxRangesPerPass = kMaxRangesPerPass; + setConfigParams.numNestingLevels = 1; + setConfigParams.minNestingLevel = 1; + CUPTI_API_CALL(cuptiRangeProfilerSetConfig(&setConfigParams)); + } + + static void EndKernelMetrics() { + ZoneScoped; + auto& rangeProfiler = PersistentState::Get().rangeProfiler; + if (rangeProfiler == nullptr) { + return; + } + // TODO(metrics): cuptiRangeProfilerStop() once Start() is wired up. + CUpti_RangeProfiler_Disable_Params destroyParams = { CUpti_RangeProfiler_Disable_Params_STRUCT_SIZE }; + destroyParams.pRangeProfilerObject = rangeProfiler; + CUPTI_API_CALL(cuptiRangeProfilerDisable(&destroyParams)); + rangeProfiler = nullptr; + + auto& profilerHostObject = PersistentState::Get().profilerHostObject; + if (profilerHostObject != nullptr) { + CUpti_Profiler_Host_Deinitialize_Params hostDeinitParams = { CUpti_Profiler_Host_Deinitialize_Params_STRUCT_SIZE }; + hostDeinitParams.pHostObject = profilerHostObject; + CUPTI_API_CALL(cuptiProfilerHostDeinitialize(&hostDeinitParams)); + profilerHostObject = nullptr; + } + } + + // Debug utility: print host/chip information and the metrics available + // for collection on the current device. Self-contained — spins up a + // temporary host object and tears it down, so it is safe to call without + // an active profiling session. Output goes to stdout and can be verbose. + static void DumpAvailableMetrics() { + ZoneScoped; + CUcontext cuCtx = nullptr; + DRIVER_API_CALL(cuCtxGetCurrent(&cuCtx)); + if (cuCtx == nullptr) { + fprintf(stderr, "TracyCUDA: no current CUDA context; cannot list metrics.\n"); + return; + } + + // Device identity (human-readable), independent of the profiler. + CUdevice cuDevice = 0; + DRIVER_API_CALL(cuCtxGetDevice(&cuDevice)); + char deviceName[256] = {}; + DRIVER_API_CALL(cuDeviceGetName(deviceName, sizeof(deviceName), cuDevice)); + int ccMajor = 0, ccMinor = 0; + DRIVER_API_CALL(cuDeviceGetAttribute(&ccMajor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevice)); + DRIVER_API_CALL(cuDeviceGetAttribute(&ccMinor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevice)); + + // Profiler-side chip name (what the host object keys metrics on). + CUpti_Profiler_Initialize_Params initParams = { CUpti_Profiler_Initialize_Params_STRUCT_SIZE }; + CUPTI_API_CALL(cuptiProfilerInitialize(&initParams)); + CUpti_Device_GetChipName_Params chipNameParams = { CUpti_Device_GetChipName_Params_STRUCT_SIZE }; + chipNameParams.deviceIndex = (size_t)cuDevice; + CUPTI_API_CALL(cuptiDeviceGetChipName(&chipNameParams)); + + fprintf(stdout, "\nTracyCUDA kernel-metrics host info:\n"); + fprintf(stdout, " Device : %s (compute capability %d.%d)\n", deviceName, ccMajor, ccMinor); + fprintf(stdout, " Chip : %s\n", chipNameParams.pChipName); + + // Listing metrics is a pure host-side query keyed on the chip name, + // so it must NOT touch the device: we deliberately skip + // cuptiProfilerGetCounterAvailability() (which reserves the perf + // counters and is admin-gated -> INSUFFICIENT_PRIVILEGES / HARDWARE_BUSY) + // and pass a null counter-availability image. Per the CUPTI docs the + // image is only required for chips newer than the CUPTI build; for a + // chip known at release time the name alone enumerates every metric. + // The listing therefore reflects everything the chip defines rather + // than what this context can currently collect — which is exactly + // what a static "available metrics" dump should report. + CUpti_Profiler_Host_Initialize_Params hostInitParams = { CUpti_Profiler_Host_Initialize_Params_STRUCT_SIZE }; + hostInitParams.profilerType = CUPTI_PROFILER_TYPE_RANGE_PROFILER; + hostInitParams.pChipName = chipNameParams.pChipName; + hostInitParams.pCounterAvailabilityImage = nullptr; + CUPTI_API_CALL(cuptiProfilerHostInitialize(&hostInitParams)); + auto* hostObject = hostInitParams.pHostObject; + + // Enumerate base metrics per type, and each base metric's submetrics + // (the fully-qualified, collectible names like ".avg.pct_..."). + const CUpti_MetricType metricTypes[] = { + CUPTI_METRIC_TYPE_COUNTER, + CUPTI_METRIC_TYPE_RATIO, + CUPTI_METRIC_TYPE_THROUGHPUT, + }; + const char* const metricTypeNames[] = { "counter", "ratio", "throughput" }; + for (size_t t = 0; t < sizeof(metricTypes) / sizeof(metricTypes[0]); ++t) { + CUpti_Profiler_Host_GetBaseMetrics_Params baseParams = { CUpti_Profiler_Host_GetBaseMetrics_Params_STRUCT_SIZE }; + baseParams.pHostObject = hostObject; + baseParams.metricType = metricTypes[t]; + CUPTI_API_CALL(cuptiProfilerHostGetBaseMetrics(&baseParams)); + fprintf(stdout, "\n %zu %s metrics:\n", (size_t)baseParams.numMetrics, metricTypeNames[t]); + for (size_t i = 0; i < baseParams.numMetrics; ++i) { + const char* baseName = baseParams.ppMetricNames[i]; + fprintf(stdout, " %s\n", baseName); + CUpti_Profiler_Host_GetSubMetrics_Params subParams = { CUpti_Profiler_Host_GetSubMetrics_Params_STRUCT_SIZE }; + subParams.pHostObject = hostObject; + subParams.metricType = metricTypes[t]; + subParams.pMetricName = baseName; + CUPTI_API_CALL(cuptiProfilerHostGetSubMetrics(&subParams)); + for (size_t s = 0; s < subParams.numOfSubmetrics; ++s) { + fprintf(stdout, " .%s\n", subParams.ppSubMetrics[s]); + } + } + } + + CUpti_Profiler_Host_Deinitialize_Params hostDeinitParams = { CUpti_Profiler_Host_Deinitialize_Params_STRUCT_SIZE }; + hostDeinitParams.pHostObject = hostObject; + CUPTI_API_CALL(cuptiProfilerHostDeinitialize(&hostDeinitParams)); + } + #endif//TRACY_CUDA_ENABLE_KERNEL_METRICS + static void BeginInstrumentation(CUDACtx* profilerHost) { auto& currentProfilerHost = PersistentState::Get().profilerHost; if (currentProfilerHost != nullptr) { @@ -1274,6 +1568,8 @@ namespace tracy // CUDA API calls and device activities that happens past this point cudaDeviceSynchronize(); + DumpAvailableMetrics(); + auto& subscriber = PersistentState::Get().subscriber; CUPTI_API_CALL(cuptiSubscribe(&subscriber, CUPTI::OnCallbackAPI, profilerHost)); CUPTI_API_CALL(cuptiActivityRegisterCallbacks(CUPTI::OnBufferRequested, CUPTI::OnBufferCompleted)); @@ -1399,6 +1695,25 @@ namespace tracy CUpti_SubscriberHandle subscriber = {}; CUDACtx* profilerHost = nullptr; + #if TRACY_CUDA_ENABLE_KERNEL_METRICS + // CUPTI Range Profiler object, created in BeginKernelMetrics() and + // destroyed in EndKernelMetrics(). nullptr when metrics are inactive. + CUpti_RangeProfiler_Object* rangeProfiler = nullptr; + // Host-side profiler object: builds metric config images from a + // metric list and decodes counter data into metric values. It is + // chip-specific, hence the cached chip name and counter availability + // snapshot it is initialized from. + CUpti_Profiler_Host_Object* profilerHostObject = nullptr; + std::string chipName; + std::vector counterAvailabilityImage; + // Config image: which counters to collect (baked from the metric + // list). Counter-data image: where the Range Profiler writes the + // collected values. Both must outlive the active profiling session, + // as cuptiRangeProfilerSetConfig() retains pointers into them. + std::vector configImage; + std::vector counterDataImage; + #endif//TRACY_CUDA_ENABLE_KERNEL_METRICS + Collector collector; static PersistentState& Get() { @@ -1491,6 +1806,10 @@ namespace tracy #define TracyCUDAStartProfiling(ctx) ctx->StartProfiling() #define TracyCUDAStopProfiling(ctx) ctx->StopProfiling() +#define TracyCUDAEnableKernelMetrics(ctx) ctx->EnableKernelMetrics() +#define TracyCUDADisableKernelMetrics(ctx) ctx->DisableKernelMetrics() +#define TracyCUDADumpKernelMetrics(ctx) ctx->DumpKernelMetrics() + #define TracyCUDACollect(ctx) ctx->Collect() #endif