CUPTI PC 采样教程

GitHub 仓库和完整教程可在 https://github.com/eunomia-bpf/cupti-tutorial 获取。

简介

在优化CUDA内核时，了解代码的哪些部分消耗最多执行时间至关重要。程序计数器（PC）采样是一种强大的技术，允许您以最小的性能开销识别GPU代码中的热点。本教程演示如何使用CUPTI的PC采样API在内核执行期间收集程序计数器样本并将其映射回源代码，帮助您将优化工作集中在最有影响的地方。

您将学到什么

如何为CUDA内核配置和启用PC采样
在内核执行期间收集和处理PC样本
将PC地址映射到源代码位置
识别CUDA代码中的性能热点
分析采样数据以指导优化工作

理解PC采样

程序计数器（PC）采样通过定期记录每个活跃线程当前执行的指令来工作。此过程如下进行：

GPU硬件定期采样活跃warp的程序计数器
这些样本在内核执行期间被收集和缓冲
内核完成后，分析样本以确定哪些代码区域被执行得最频繁
使用调试信息将PC地址映射回源代码

与基于仪器的性能分析不同，PC采样对内核性能的影响最小，并提供GPU执行时间消耗的统计视图。

代码演练

1. 设置PC采样

首先，我们需要配置和启用PC采样：

CUpti_PCSamplingConfig config;
memset(&config, 0, sizeof(config));
config.size = sizeof(config);
config.samplingPeriod = 5;  // 每5个周期采样一次
config.samplingPeriod2 = 0; // 此示例中未使用
config.samplingPeriodRatio = 0; // 此示例中未使用
config.collectingMode = CUPTI_PC_SAMPLING_COLLECTION_MODE_KERNEL;
config.samplingBufferSize = 0x2000; // 8KB缓冲区
config.stoppingCount = 0; // 无停止条件

// 获取当前CUDA上下文
CUcontext context;
DRIVER_API_CALL(cuCtxGetCurrent(&context));

// 启用PC采样
CUPTI_CALL(cuptiPCSamplingEnable(&config));

此代码： 1. 创建PC采样的配置结构 2. 设置采样周期（多久采样一次PC） 3. 配置收集模式以在内核执行期间采样 4. 设置存储样本的缓冲区大小 5. 在当前CUDA上下文上启用PC采样

2. 注册缓冲区回调

要处理PC采样数据，我们需要注册在缓冲区满时将被调用的回调：

// 注册缓冲区处理的回调
CUPTI_CALL(cuptiPCSamplingRegisterBufferHandler(handleBuffer, userData));

// 处理PC采样缓冲区的回调函数
void handleBuffer(uint8_t *buffer, size_t size, size_t validSize, void *userData)
{
    PCData *pcData = (PCData *)userData;

    // 处理缓冲区中的所有记录
    CUpti_PCSamplingData *record = (CUpti_PCSamplingData *)buffer;

    while (validSize >= sizeof(CUpti_PCSamplingData)) {
        // 处理PC样本
        processPCSample(record, pcData);

        // 移动到下一个记录
        validSize -= sizeof(CUpti_PCSamplingData);
        record++;
    }
}

此代码： 1. 注册一个在PC采样缓冲区准备就绪时将被调用的回调函数 2. 在回调中，处理每个PC采样记录 3. 根据收集的样本更新统计信息

3. 处理PC样本

对于每个PC样本，我们需要处理和存储信息：

void processPCSample(CUpti_PCSamplingData *sample, PCData *pcData)
{
    // 从样本中提取信息
    uint64_t pc = sample->pc;
    uint32_t functionId = sample->functionId;
    uint32_t stall = sample->stallReason;

    // 更新PC直方图
    pcData->totalSamples++;

    // 为此PC查找或创建条目
    PCEntry *entry = findPCEntry(pcData, pc);
    if (!entry) {
        entry = createPCEntry(pcData, pc);
    }

    // 更新此PC的样本计数
    entry->sampleCount++;

    // 如果需要，更新阻塞原因计数
    if (stall != CUPTI_PC_SAMPLING_STALL_NONE) {
        entry->stallCount[stall]++;
    }
}

此函数： 1. 从样本中提取程序计数器（PC）值 2. 更新总样本计数 3. 在我们的数据结构中为此PC查找或创建条目 4. 更新此PC的统计信息，包括阻塞原因

4. 将PC映射到源代码

要使PC值有意义，我们需要将它们映射回源代码：

void mapPCsToSource(PCData *pcData)
{
    // 获取内核的CUDA模块
    CUmodule module;
    DRIVER_API_CALL(cuModuleGetFunction(&module, pcData->function));

    // 对于每个PC条目，获取源信息
    for (int i = 0; i < pcData->numEntries; i++) {
        PCEntry *entry = &pcData->entries[i];

        // 获取源文件和行信息
        CUpti_LineInfo lineInfo;
        lineInfo.size = sizeof(lineInfo);

        CUPTI_CALL(cuptiGetLineInfo(module, entry->pc, &lineInfo));

        // 存储源信息
        if (lineInfo.lineInfoValid) {
            entry->fileName = strdup(lineInfo.fileName);
            entry->lineNumber = lineInfo.lineNumber;
            entry->functionName = strdup(lineInfo.functionName);
        } else {
            entry->fileName = strdup("unknown");
            entry->lineNumber = 0;
            entry->functionName = strdup("unknown");
        }
    }
}

此函数： 1. 获取我们内核的CUDA模块 2. 对于每个PC条目，调用CUPTI获取源行信息 3. 为每个PC存储文件名、行号和函数名

5. 带热点的示例内核

为了演示PC采样，我们将使用一个带有一些故意热点的内核：

__global__ void sampleKernel(float *data, int size, int iterations)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx < size) {
        float value = data[idx];

        // 热点1：计算密集型循环
        for (int i = 0; i < iterations; i++) {
            value = value * value + 0.5f;
        }

        // 热点2：带发散的条件分支
        if (idx % 2 == 0) {
            for (int i = 0; i < iterations / 2; i++) {
                value = sinf(value);
            }
        } else {
            for (int i = 0; i < iterations / 4; i++) {
                value = cosf(value);
            }
        }

        // 热点3：内存访问模式
        int offset = (idx * 17) % size;
        value += data[offset];

        data[idx] = value;
    }
}

此内核具有几个将在PC采样中显示的特征： 1. 一个将成为主要热点的计算密集型循环 2. 导致线程发散的条件分支 3. 非合并内存访问模式

6. 分析和显示结果

收集样本后，我们分析和显示结果：

void analyzeResults(PCData *pcData)
{
    // 按样本计数排序PC条目
    qsort(pcData->entries, pcData->numEntries, sizeof(PCEntry), comparePCEntries);

    // 打印摘要
    printf("\nPC采样结果：\n");
    printf("  收集的总样本数：%llu\n", pcData->totalSamples);
    printf("  唯一PC地址：%d\n\n", pcData->numEntries);

    // 打印前5个热点
    printf("前5个热点：\n");
    int numToShow = (pcData->numEntries < 5) ? pcData->numEntries : 5;

    for (int i = 0; i < numToShow; i++) {
        PCEntry *entry = &pcData->entries[i];
        float percentage = 100.0f * entry->sampleCount / pcData->totalSamples;

        printf("  %d. PC: 0x%llx（%.1f%%的样本）- %s:%d - %s\n",
               i + 1, entry->pc, percentage,
               entry->fileName, entry->lineNumber, entry->functionName);

        // 如果可用，打印阻塞原因
        if (entry->stallCount[CUPTI_PC_SAMPLING_STALL_MEMORY_THROTTLE] > 0) {
            printf("     内存限制阻塞：%d\n", 
                  entry->stallCount[CUPTI_PC_SAMPLING_STALL_MEMORY_THROTTLE]);
        }
        if (entry->stallCount[CUPTI_PC_SAMPLING_STALL_SYNC] > 0) {
            printf("     同步阻塞：%d\n", 
                  entry->stallCount[CUPTI_PC_SAMPLING_STALL_SYNC]);
        }
        // ... 打印其他阻塞原因 ...
    }
}

Share on Share on