CUPTI 用户范围分析教程
GitHub 仓库和完整教程可在 https://github.com/eunomia-bpf/cupti-tutorial 获取。
简介
在分析 CUDA 应用程序时,您通常需要专注于代码的特定部分,而不是整个内核或完整的应用程序。CUPTI 的用户范围分析功能允许您在代码中定义自定义范围,并仅在这些范围内收集性能指标。这为您提供了对应用程序哪些部分被分析的精确控制,使得在复杂应用程序中更容易识别和优化性能瓶颈。本教程演示如何使用 CUPTI 的分析器 API 定义自定义范围并在其中收集性能指标。
您将学到什么
- 如何在 CUDA 代码中定义和仪器化用户指定的范围
- 为用户范围分析设置 CUPTI 分析器 API
- 在您定义的范围内收集 GPU 性能指标
- 处理和分析每个范围收集的指标
- 比较应用程序不同部分的性能
理解用户范围分析
与自动分析单个 CUDA API 调用或内核不同,用户范围分析让您:
- 定义可能包含多个 CUDA 操作的代码逻辑段
- 为这些段提供有意义的名称以便于分析
- 将分析资源集中在应用程序最重要的部分
- 比较不同算法方法的性能指标
这种方法在以下情况下特别有用: - 您的应用程序具有不同性能特征的不同阶段 - 您想比较同一算法的不同实现 - 您需要将特定的操作序列作为单个单元进行分析 - 您正在优化大型应用程序的特定部分
代码演练
1. 设置 CUPTI 分析器
首先,我们需要初始化 CUPTI 分析器 API 并为用户范围分析配置它:
CUpti_Profiler_Initialize_Params initializeParams = {CUpti_Profiler_Initialize_Params_STRUCT_SIZE};
CUPTI_API_CALL(cuptiProfilerInitialize(&initializeParams));
// 获取当前设备的芯片名称
CUpti_Device_GetChipName_Params getChipNameParams = {CUpti_Device_GetChipName_Params_STRUCT_SIZE};
getChipNameParams.deviceIndex = 0;
CUPTI_API_CALL(cuptiDeviceGetChipName(&getChipNameParams));
const char *chipName = getChipNameParams.pChipName;
// 创建指标配置
NVPW_InitializeHost_Params initializeHostParams = {NVPW_InitializeHost_Params_STRUCT_SIZE};
NVPW_API_CALL(NVPW_InitializeHost(&initializeHostParams));
// 设置要收集的指标
const char *metricNames[] = {"smsp__warps_launched.avg"};
struct MetricNameList metricList;
metricList.numMetrics = 1;
metricList.metricNames = metricNames;
// 创建计数器数据映像和配置
CUpti_Profiler_CounterDataImageOptions counterDataImageOptions = {
CUpti_Profiler_CounterDataImageOptions_STRUCT_SIZE};
counterDataImageOptions.pChipName = chipName;
counterDataImageOptions.counterDataImageSize = 0;
CUPTI_API_CALL(cuptiProfilerCounterDataImageCalculateSize(&counterDataImageOptions));
counterDataImage = (uint8_t *)malloc(counterDataImageOptions.counterDataImageSize);
counterDataImageOptions.pCounterDataImage = counterDataImage;
CUPTI_API_CALL(cuptiProfilerCounterDataImageInitialize(&counterDataImageOptions));
// 为用户范围分析配置分析器
CUpti_Profiler_UserRange_Config_Params configParams = {
CUpti_Profiler_UserRange_Config_Params_STRUCT_SIZE};
configParams.pCounterDataPrefixImage = counterDataImage;
configParams.counterDataPrefixImageSize = counterDataImageOptions.counterDataImageSize;
configParams.maxRangesPerPass = 1;
configParams.maxLaunchesPerPass = 1;
CUPTI_API_CALL(cuptiProfilerUserRangeConfigureScratchBuffer(&configParams));
这段代码: 1. 初始化 CUPTI 分析器 API 2. 获取当前设备的芯片名称 3. 设置我们想要收集的指标(在这种情况下是 "smsp__warps_launched.avg") 4. 创建并配置计数器数据映像 5. 专门为用户范围分析配置分析器
2. 在代码中定义用户范围
接下来,我们围绕想要分析的代码段定义范围:
void profileVectorOperations(int *d_A, int *d_B, int *d_C, int numElements)
{
// 定义范围名称
const char *rangeName = "Vector Add-Subtract";
// 开始用户范围
CUpti_Profiler_BeginRange_Params beginRangeParams = {CUpti_Profiler_BeginRange_Params_STRUCT_SIZE};
beginRangeParams.pRangeName = rangeName;
CUPTI_API_CALL(cuptiProfilerBeginRange(&beginRangeParams));
// 在范围内执行操作
int threadsPerBlock = 256;
int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;
printf("启动内核:块数 %d,每块线程数 %d\n", blocksPerGrid, threadsPerBlock);
// 第一个向量操作:加法
VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);
// 第二个向量操作:减法
VecSub<<<blocksPerGrid, threadsPerBlock>>>(d_C, d_B, d_C, numElements);
// 结束用户范围
CUpti_Profiler_EndRange_Params endRangeParams = {CUpti_Profiler_EndRange_Params_STRUCT_SIZE};
CUPTI_API_CALL(cuptiProfilerEndRange(&endRangeParams));
}
这个函数:
1. 为范围定义有意义的名称("Vector Add-Subtract")
2. 使用 cuptiProfilerBeginRange
开始用户范围
3. 在范围内执行多个 GPU 操作(两个内核启动)
4. 使用 cuptiProfilerEndRange
结束范围
3. 启动和停止分析器会话
我们需要在范围之前启动分析器,之后停止它:
int main(int argc, char *argv[])
{
// 初始化 CUDA 和 CUPTI
initializeCuda();
initializeProfiler();
// 分配和初始化数据
int *h_A, *h_B, *h_C;
int *d_A, *d_B, *d_C;
// ... 分配和初始化代码 ...
// 开始分析器会话
CUpti_Profiler_BeginSession_Params beginSessionParams = {CUpti_Profiler_BeginSession_Params_STRUCT_SIZE};
beginSessionParams.counterDataImageSize = counterDataImageOptions.counterDataImageSize;
beginSessionParams.pCounterDataImage = counterDataImage;
beginSessionParams.maxRangesPerPass = 1;
beginSessionParams.maxLaunchesPerPass = 1;
CUPTI_API_CALL(cuptiProfilerBeginSession(&beginSessionParams));
// 启用分析
CUpti_Profiler_EnableProfiling_Params enableProfilingParams = {
CUpti_Profiler_EnableProfiling_Params_STRUCT_SIZE};
CUPTI_API_CALL(cuptiProfilerEnableProfiling(&enableProfilingParams));
// 执行包含用户范围的分析代码
profileVectorOperations(d_A, d_B, d_C, numElements);
// 禁用分析
CUpti_Profiler_DisableProfiling_Params disableProfilingParams = {
CUpti_Profiler_DisableProfiling_Params_STRUCT_SIZE};
CUPTI_API_CALL(cuptiProfilerDisableProfiling(&disableProfilingParams));
// 结束分析器会话
CUpti_Profiler_EndSession_Params endSessionParams = {CUpti_Profiler_EndSession_Params_STRUCT_SIZE};
CUPTI_API_CALL(cuptiProfilerEndSession(&endSessionParams));
// 处理结果
processCounterData();
// 清理
// ... 清理代码 ...
return 0;
}
这段代码:
1. 初始化 CUDA 和 CUPTI 分析器
2. 使用 cuptiProfilerBeginSession
开始分析器会话
3. 使用 cuptiProfilerEnableProfiling
启用分析
4. 执行包含我们用户范围的代码
5. 禁用分析并结束会话
6. 处理收集的数据
4. 处理收集的指标
分析后,我们需要处理计数器数据以计算指标:
void processCounterData()
{
// 获取当前设备的芯片名称
CUpti_Device_GetChipName_Params getChipNameParams = {CUpti_Device_GetChipName_Params_STRUCT_SIZE};
getChipNameParams.deviceIndex = 0;
CUPTI_API_CALL(cuptiDeviceGetChipName(&getChipNameParams));
const char *chipName = getChipNameParams.pChipName;
// 初始化评估上下文
NVPW_CUDA_MetricsEvaluator_CalculateScratchBufferSize_Params scratchBufferSizeParams = {
NVPW_CUDA_MetricsEvaluator_CalculateScratchBufferSize_Params_STRUCT_SIZE};
scratchBufferSizeParams.pChipName = chipName;
scratchBufferSizeParams.pCounterDataImage = counterDataImage;
NVPW_API_CALL(NVPW_CUDA_MetricsEvaluator_CalculateScratchBufferSize(&scratchBufferSizeParams));
// 分配和初始化评估器
void *scratchBuffer = malloc(scratchBufferSizeParams.scratchBufferSize);
NVPW_CUDA_MetricsEvaluator_Initialize_Params initializeParams = {
NVPW_CUDA_MetricsEvaluator_Initialize_Params_STRUCT_SIZE};
initializeParams.scratchBufferSize = scratchBufferSizeParams.scratchBufferSize;
initializeParams.pScratchBuffer = scratchBuffer;
initializeParams.pChipName = chipName;
initializeParams.pCounterDataImage = counterDataImage;
NVPW_API_CALL(NVPW_CUDA_MetricsEvaluator_Initialize(&initializeParams));
// 评估指标
NVPW_MetricsEvaluator_EvaluateToGpuValues_Params evaluateParams = {
NVPW_MetricsEvaluator_EvaluateToGpuValues_Params_STRUCT_SIZE};
evaluateParams.pMetricNames = metricNames;
evaluateParams.numMetrics = 1;
evaluateParams.pCounterDataImage = counterDataImage;
NVPW_API_CALL(NVPW_MetricsEvaluator_EvaluateToGpuValues(&evaluateParams));
// 打印结果
printf("用户范围分析结果:\n");
for (size_t i = 0; i < evaluateParams.numMetrics; i++) {
printf(" %s: %.2f\n", metricNames[i], evaluateParams.pGpuValues[i]);
}
// 清理
free(scratchBuffer);
}
这段代码: 1. 获取设备芯片名称以进行指标评估 2. 计算并分配评估器的暂存缓冲区 3. 初始化指标评估器 4. 评估我们指定的指标 5. 打印结果并清理资源
高级用例
多范围分析
class MultiRangeProfiler {
private:
std::vector<std::string> rangeNames;
std::map<std::string, std::map<std::string, double>> rangeMetrics;
public:
void profileAlgorithmComparison() {
// 定义要分析的算法变体
std::vector<std::string> algorithms = {
"Naive Implementation",
"Shared Memory Optimized",
"Coalesced Access",
"Tiled Algorithm"
};
for (const auto& algorithm : algorithms) {
// 开始范围
CUpti_Profiler_BeginRange_Params beginParams = {
CUpti_Profiler_BeginRange_Params_STRUCT_SIZE};
beginParams.pRangeName = algorithm.c_str();
CUPTI_API_CALL(cuptiProfilerBeginRange(&beginParams));
// 执行特定的算法实现
if (algorithm == "Naive Implementation") {
runNaiveKernel();
} else if (algorithm == "Shared Memory Optimized") {
runSharedMemoryKernel();
} else if (algorithm == "Coalesced Access") {
runCoalescedKernel();
} else if (algorithm == "Tiled Algorithm") {
runTiledKernel();
}
// 结束范围
CUpti_Profiler_EndRange_Params endParams = {
CUpti_Profiler_EndRange_Params_STRUCT_SIZE};
CUPTI_API_CALL(cuptiProfilerEndRange(&endParams));
// 收集此范围的指标
auto metrics = collectRangeMetrics();
rangeMetrics[algorithm] = metrics;
}
// 比较结果
compareAlgorithmPerformance();
}
private:
void compareAlgorithmPerformance() {
printf("=== 算法性能比较 ===\n");
// 找到最佳性能作为基线
double bestTime = std::numeric_limits<double>::max();
std::string bestAlgorithm;
for (const auto& [name, metrics] : rangeMetrics) {
double execTime = metrics.at("sm__cycles_elapsed.avg");
if (execTime < bestTime) {
bestTime = execTime;
bestAlgorithm = name;
}
}
// 打印相对性能
for (const auto& [name, metrics] : rangeMetrics) {
double execTime = metrics.at("sm__cycles_elapsed.avg");
double speedup = bestTime / execTime;
printf("%s:\n", name.c_str());
printf(" 执行时间: %.2f 周期\n", execTime);
printf(" 相对于最佳的速度提升: %.2fx", speedup);
if (name == bestAlgorithm) {
printf(" (最佳)");
}
printf("\n\n");
}
}
};
性能回归检测
class PerformanceRegressionDetector {
private:
std::map<std::string, double> baselineMetrics;
double regressionThreshold = 0.95; // 5% 性能下降阈值
public:
void setBaseline(const std::string& rangeName) {
// 运行基线测试
profileRange(rangeName);
auto metrics = collectRangeMetrics();
baselineMetrics[rangeName] = metrics["sm__cycles_elapsed.avg"];
printf("为 %s 设置基线: %.2f 周期\n",
rangeName.c_str(), baselineMetrics[rangeName]);
}
bool checkForRegression(const std::string& rangeName) {
// 运行当前版本
profileRange(rangeName);
auto metrics = collectRangeMetrics();
double currentPerformance = metrics["sm__cycles_elapsed.avg"];
// 比较与基线(较低的周期 = 更好的性能)
double baseline = baselineMetrics[rangeName];
double performanceRatio = baseline / currentPerformance;
printf("性能检查 %s:\n", rangeName.c_str());
printf(" 基线: %.2f 周期\n", baseline);
printf(" 当前: %.2f 周期\n", currentPerformance);
printf(" 比率: %.3f", performanceRatio);
if (performanceRatio < regressionThreshold) {
printf(" - 检测到回归!\n");
return true;
} else {
printf(" - 通过\n");
return false;
}
}
private:
void profileRange(const std::string& rangeName) {
CUpti_Profiler_BeginRange_Params beginParams = {
CUpti_Profiler_BeginRange_Params_STRUCT_SIZE};
beginParams.pRangeName = rangeName.c_str();
CUPTI_API_CALL(cuptiProfilerBeginRange(&beginParams));
// 执行要测试的代码
executeTestCode(rangeName);
CUpti_Profiler_EndRange_Params endParams = {
CUpti_Profiler_EndRange_Params_STRUCT_SIZE};
CUPTI_API_CALL(cuptiProfilerEndRange(&endParams));
}
};
自适应优化
class AdaptiveOptimizer {
public:
void optimizeKernelLaunchParameters() {
std::vector<int> blockSizes = {128, 256, 512, 1024};
std::map<int, double> performanceResults;
for (int blockSize : blockSizes) {
std::string rangeName = "BlockSize_" + std::to_string(blockSize);
// 分析此配置
CUpti_Profiler_BeginRange_Params beginParams = {
CUpti_Profiler_BeginRange_Params_STRUCT_SIZE};
beginParams.pRangeName = rangeName.c_str();
CUPTI_API_CALL(cuptiProfilerBeginRange(&beginParams));
// 使用当前块大小运行内核
runKernelWithBlockSize(blockSize);
CUpti_Profiler_EndRange_Params endParams = {
CUpti_Profiler_EndRange_Params_STRUCT_SIZE};
CUPTI_API_CALL(cuptiProfilerEndRange(&endParams));
// 收集性能数据
auto metrics = collectRangeMetrics();
performanceResults[blockSize] = metrics["sm__cycles_elapsed.avg"];
}
// 找到最佳配置
int bestBlockSize = findOptimalBlockSize(performanceResults);
printf("最佳块大小: %d\n", bestBlockSize);
// 应用最佳配置
applyOptimalConfiguration(bestBlockSize);
}
private:
int findOptimalBlockSize(const std::map<int, double>& results) {
int bestSize = 0;
double bestTime = std::numeric_limits<double>::max();
for (const auto& [size, time] : results) {
printf("块大小 %d: %.2f 周期\n", size, time);
if (time < bestTime) {
bestTime = time;
bestSize = size;
}
}
return bestSize;
}
};
最佳实践
范围命名约定
// 使用分层命名以便于分析
class RangeNamingConventions {
public:
static std::string createRangeName(const std::string& module,
const std::string& function,
const std::string& variant = "") {
std::string name = module + "::" + function;
if (!variant.empty()) {
name += "_" + variant;
}
return name;
}
// 示例用法
void demonstrateNaming() {
// 模块级范围
profileRange("LinearAlgebra::MatrixMultiply");
profileRange("ImageProcessing::Convolution");
// 变体比较
profileRange("GEMM::Naive");
profileRange("GEMM::Tiled");
profileRange("GEMM::SharedMemory");
// 参数化测试
profileRange("Sort::QuickSort_Size1M");
profileRange("Sort::QuickSort_Size10M");
}
};
错误处理和验证
class ProfilerValidator {
public:
static bool validateProfilerState() {
// 检查分析器是否正确初始化
CUpti_Profiler_GetInitializationStatus_Params statusParams = {
CUpti_Profiler_GetInitializationStatus_Params_STRUCT_SIZE};
CUptiResult result = cuptiProfilerGetInitializationStatus(&statusParams);
if (result != CUPTI_SUCCESS) {
printf("分析器未正确初始化\n");
return false;
}
return true;
}
static void handleProfilerError(CUptiResult result, const char* operation) {
if (result != CUPTI_SUCCESS) {
const char* errorString;
cuptiGetResultString(result, &errorString);
printf("分析器错误在 %s: %s\n", operation, errorString);
exit(1);
}
}
};
故障排除
常见问题
- 范围嵌套:确保范围正确嵌套,避免重叠
- 指标兼容性:验证指标在目标 GPU 架构上可用
- 内存限制:监控分析器内存使用情况
- 性能开销:平衡分析粒度与开销
调试技巧
void debugProfilingSession() {
// 启用详细的 CUPTI 日志记录
setenv("CUPTI_DEBUG", "1", 1);
// 检查可用指标
listAvailableMetrics();
// 验证范围边界
validateRangeBoundaries();
// 监控内存使用
checkMemoryUsage();
}
用户范围分析为精确的 CUDA 性能分析提供了强大的工具。通过仔细定义有意义的范围并收集相关指标,您可以获得对应用程序性能特征的深入洞察,并做出数据驱动的优化决策。