CUPTI 回调事件教程

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

简介

跟踪 GPU 性能事件对于优化 CUDA 应用程序至关重要。本教程演示了如何使用 CUPTI 的回调机制在内核执行期间收集特定的性能事件。我们将专注于"已执行指令"指标，以展示您的内核运行了多少 GPU 指令。

您将学到的内容

如何为 CUDA 运行时 API 调用设置 CUPTI 回调函数
创建和管理用于收集性能数据的事件组
围绕内核启动收集事件数据的技术
解释事件值以进行性能分析

代码演练

1. 事件收集结构

示例设置了两个关键数据结构：

// 存储事件组和事件 ID
typedef struct cupti_eventData_st {
  CUpti_EventGroup eventGroup;
  CUpti_EventID eventId;
} cupti_eventData;

// 存储回调收集的事件数据和值
typedef struct RuntimeApiTrace_st {
  cupti_eventData *eventData;
  uint64_t eventVal;
} RuntimeApiTrace_t;

这些结构在整个应用程序中维护事件状态和收集的值。

2. 回调函数

示例的核心是回调函数，当调用某些 CUDA 运行时 API 函数时会被调用：

void CUPTIAPI getEventValueCallback(void *userdata, CUpti_CallbackDomain domain,
                                    CUpti_CallbackId cbid, const CUpti_CallbackData *cbInfo)
{
  // 仅处理内核启动的回调
  if ((cbid != CUPTI_RUNTIME_TRACE_CBID_cudaLaunch_v3020) &&
      (cbid != CUPTI_RUNTIME_TRACE_CBID_cudaLaunchKernel_v7000)) {
    return;
  }

  RuntimeApiTrace_t *traceData = (RuntimeApiTrace_t*)userdata;

  // 进入 CUDA 运行时函数时（内核启动前）
  if (cbInfo->callbackSite == CUPTI_API_ENTER) {
    // 同步设备以确保干净的事件收集
    cudaDeviceSynchronize();

    // 将收集模式设置为内核级别
    cuptiSetEventCollectionMode(cbInfo->context, CUPTI_EVENT_COLLECTION_MODE_KERNEL);

    // 启用事件组开始收集数据
    cuptiEventGroupEnable(traceData->eventData->eventGroup);
  }

  // 退出 CUDA 运行时函数时（内核完成后）
  if (cbInfo->callbackSite == CUPTI_API_EXIT) {
    // 确定发生了多少个事件实例
    uint32_t numInstances = 0;
    cuptiEventGroupGetAttribute(traceData->eventData->eventGroup, 
                               CUPTI_EVENT_GROUP_ATTR_INSTANCE_COUNT, 
                               &valueSize, &numInstances);

    // 为事件值分配空间
    uint64_t *values = (uint64_t *) malloc(sizeof(uint64_t) * numInstances);

    // 确保内核完成
    cudaDeviceSynchronize();

    // 读取事件值
    cuptiEventGroupReadEvent(traceData->eventData->eventGroup, 
                            CUPTI_EVENT_READ_FLAG_NONE, 
                            traceData->eventData->eventId, 
                            &bytesRead, values);

    // 聚合所有实例的值
    traceData->eventVal = 0;
    for (i=0; i<numInstances; i++) {
      traceData->eventVal += values[i];
    }

    // 清理
    free(values);
    cuptiEventGroupDisable(traceData->eventData->eventGroup);
  }
}

该函数执行以下关键操作： - 在 API 进入时（内核启动前）：启用事件收集 - 在 API 退出时（内核完成后）：读取事件值并禁用收集

3. 设置事件和回调

int main(int argc, char *argv[])
{
  // 要跟踪的默认事件（已执行指令）
  const char *eventName = EVENT_NAME;

  // 解析设备和事件的命令行参数
  if (argc > 2)
    eventName = argv[2];

  // 设置 CUPTI 订阅者
  cuptiSubscribe(&subscriber, (CUpti_CallbackFunc)getEventValueCallback, &trace);

  // 为内核启动启用回调
  cuptiEnableDomain(1, subscriber, CUPTI_CB_DOMAIN_RUNTIME_API);

  // 查找事件并创建事件组
  cuptiEventGetIdFromName(device, eventName, &cuptiEvent.eventId);
  cuptiEventGroupCreate(context, &cuptiEvent.eventGroup, 0);
  cuptiEventGroupAddEvent(cuptiEvent.eventGroup, cuptiEvent.eventId);

  // 将事件数据存储在我们的跟踪结构中
  trace.eventData = &cuptiEvent;
  trace.eventVal = 0;

  // 运行向量加法内核
  // ...

  // 显示事件值
  displayEventVal(&trace, eventName);

  // 清理
  cuptiEventGroupDestroy(cuptiEvent.eventGroup);
  cuptiUnsubscribe(subscriber);

  return 0;
}

主函数设置事件收集系统，运行内核，并显示结果。

4. 测试内核

示例使用简单的向量加法内核来演示事件收集：

__global__ void VecAdd(const int* A, const int* B, int* C, int N)
{
  int i = blockDim.x * blockIdx.x + threadIdx.x;
  if (i < N)
    C[i] = A[i] + B[i];
}

运行示例

构建示例：
```
make
```
使用默认参数运行（设备 0，跟踪 "inst_executed"）：
```
./callback_event
```

您可以指定不同的设备或事件：

./callback_event [device_num] [event_name]

例如：

./callback_event 0 inst_executed_global_loads

理解输出

示例产生类似于以下的输出：

CUDA Device Number: 0
CUDA Device Name: NVIDIA GeForce RTX 3080
Compute capability: 8.6
Event Name : inst_executed
Event Value : 2048000

这显示了： 1. 用于测试的设备 2. 计算能力（架构版本） 3. 被跟踪的事件（"inst_executed" = 已执行指令） 4. 内核执行的指令总数

性能见解

指令计数分析

inst_executed 事件提供了关于内核计算复杂性的见解：

高指令计数：可能表示计算密集型内核
低指令计数：可能表示内存绑定或简单操作
与理论值的比较：帮助识别编译器优化机会

事件类型

CUPTI 支持多种有用的事件类型：

// 内存相关事件
"global_load"           // 全局内存加载
"global_store"          // 全局内存存储
"shared_load"           // 共享内存加载
"shared_store"          // 共享内存存储

// 指令执行事件  
"inst_executed"         // 总执行指令
"inst_integer"          // 整数指令
"inst_fp_32"            // 32位浮点指令
"inst_fp_64"            // 64位浮点指令

// 控制流事件
"divergent_branch"      // 分支发散
"thread_inst_executed"  // 每线程指令

高级使用

多事件收集

您可以同时收集多个事件：

typedef struct MultiEventData_st {
    CUpti_EventGroup eventGroup;
    CUpti_EventID eventIds[MAX_EVENTS];
    uint64_t eventValues[MAX_EVENTS];
    int numEvents;
} MultiEventData;

void setupMultipleEvents(MultiEventData* data, const char* eventNames[], int count) {
    cuptiEventGroupCreate(context, &data->eventGroup, 0);

    for (int i = 0; i < count; i++) {
        cuptiEventGetIdFromName(device, eventNames[i], &data->eventIds[i]);
        cuptiEventGroupAddEvent(data->eventGroup, data->eventIds[i]);
    }
    data->numEvents = count;
}

事件归一化

将原始计数转换为有意义的指标：

void analyzePerformance(RuntimeApiTrace_t* trace, int numThreads) {
    double instructionsPerThread = (double)trace->eventVal / numThreads;
    double computationalIntensity = instructionsPerThread / memoryOperations;

    printf("Instructions per thread: %.2f\n", instructionsPerThread);
    printf("Computational intensity: %.2f\n", computationalIntensity);
}

时间关联

将事件数据与时间信息相结合：

typedef struct TimedEventData_st {
    uint64_t eventVal;
    uint64_t startTime;
    uint64_t endTime;
    double duration;
    double eventRate;
} TimedEventData;

void calculateEventRate(TimedEventData* data) {
    data->duration = (data->endTime - data->startTime) * 1e-9; // 转换为秒
    data->eventRate = data->eventVal / data->duration; // 每秒事件数
}

故障排除

常见问题

事件不可用：

CUptiResult result = cuptiEventGetIdFromName(device, eventName, &eventId);
if (result != CUPTI_SUCCESS) {
    printf("Event '%s' not available on this device\n", eventName);
    // 列出可用事件
    listAvailableEvents(device);
}

权限错误：
确保以管理员权限运行
检查 GPU 是否支持性能监控

计数器冲突：

// 检查事件是否可以同时收集
CUpti_EventGroupSets* eventGroupSets;
cuptiEventGroupSetsCreate(context, sizeof(eventIds), eventIds, &eventGroupSets);

if (eventGroupSets->numSets > 1) {
    printf("Events require multiple passes to collect\n");
}

调试提示

使用 cuptiGetResultString() 获取详细的错误信息
在生产代码中实现事件可用性检查
考虑不同 GPU 架构的事件兼容性

性能考虑

开销分析

事件收集会产生一些开销：

回调开销：每个 API 调用的额外处理
同步成本：cudaDeviceSynchronize() 调用
内存开销：事件值存储

优化技巧

选择性监控：

// 仅为感兴趣的内核启用回调
if (strstr(cbInfo->functionName, "TargetKernel") != NULL) {
    // 启用事件收集
}

批量处理：

// 批量收集多个启动的事件
static int launchCount = 0;
if (++launchCount % BATCH_SIZE == 0) {
    processCollectedEvents();
}

异步收集：

// 使用单独的线程处理事件数据
std::thread processingThread(processEventData, eventBuffer);
processingThread.detach();

实际应用

代码优化工作流程

基线测量：使用 inst_executed 建立基线
识别热点：找到指令计数高的内核
优化目标：专注于高影响的内核
验证改进：比较优化前后的指令计数

与其他指标的集成

将事件数据与其他性能指标结合：

typedef struct PerformanceProfile_st {
    uint64_t instructions;
    uint64_t globalLoads;
    uint64_t globalStores;
    double kernelTime;
    double ipc;  // 每周期指令数
    double bandwidth; // 内存带宽利用率
} PerformanceProfile;

总结

CUPTI 回调事件收集为了解 GPU 性能提供了详细的见解。通过跟踪特定事件如已执行指令，开发者可以：

量化计算复杂性
识别优化机会
验证性能改进
了解硬件利用率

结合其他 CUPTI 功能，事件回调形成了一个强大的性能分析工具包，帮助优化 CUDA 应用程序以获得最佳性能。

Share on Share on