CUPTI 事件采样教程
GitHub 仓库和完整教程可在 https://github.com/eunomia-bpf/cupti-tutorial 获取。
简介
在对CUDA应用程序进行性能分析时,您通常需要在应用程序运行时监控性能指标。本教程演示如何使用CUPTI的事件采样功能在内核执行期间定期收集GPU性能数据,让您实时了解应用程序的行为。
您将学到什么
- 如何在NVIDIA GPU上设置连续事件采样
- 在内核运行时监控事件的技术
- 创建多线程性能分析系统
- 解释采样事件数据以进行性能分析
理解事件采样
与在内核执行结束时给出单个值的一次性事件收集不同,事件采样允许您:
- 监控事件在内核执行期间如何随时间变化
- 检测性能变化和异常
- 将GPU活动与算法的特定阶段关联
- 观察动态工作负载的影响
代码演练
1. 设置采样线程
此示例的核心是专用采样线程,在主线程运行计算时收集事件数据:
static void *sampling_func(void *arg)
{
SamplingInfo *info = (SamplingInfo *)arg;
CUcontext context = info->context;
CUdevice device = info->device;
// 使此线程使用与主线程相同的CUDA上下文
cuCtxSetCurrent(context);
// 设置我们想要监控的事件
CUpti_EventGroup eventGroup;
CUpti_EventID eventId;
// 获取指定事件的事件ID(默认为"inst_executed")
CUPTI_CALL(cuptiEventGetIdFromName(device, info->eventName, &eventId));
// 为设备创建事件组
CUPTI_CALL(cuptiEventGroupCreate(context, &eventGroup, 0));
// 将事件添加到组中
CUPTI_CALL(cuptiEventGroupAddEvent(eventGroup, eventId));
// 设置连续收集模式(对执行期间采样至关重要)
CUPTI_CALL(cuptiEventGroupSetAttribute(eventGroup,
CUPTI_EVENT_GROUP_ATTR_COLLECTION_MODE,
sizeof(CUpti_EventCollectionMode),
&continuous));
// 启用事件组
CUPTI_CALL(cuptiEventGroupEnable(eventGroup));
// 采样直到计算完成
while (!info->terminate) {
// 读取当前事件值
size_t valueSize = sizeof(uint64_t);
uint64_t eventValue = 0;
CUPTI_CALL(cuptiEventGroupReadEvent(eventGroup,
CUPTI_EVENT_READ_FLAG_NONE,
eventId,
&valueSize,
&eventValue));
// 打印当前值
printf("%s: %llu\n", info->eventName, (unsigned long long)eventValue);
// 等待再次采样
millisleep(SAMPLE_PERIOD_MS);
}
// 清理
CUPTI_CALL(cuptiEventGroupDisable(eventGroup));
CUPTI_CALL(cuptiEventGroupDestroy(eventGroup));
return NULL;
}
此代码的关键方面:
- 上下文共享:采样线程使用与主线程相同的CUDA上下文
- 连续收集模式:启用在内核运行时读取事件值
- 定期采样:以固定间隔读取事件值(默认50ms)
- 非阻塞:采样不会中断内核执行
2. 计算线程(主线程)
主线程运行我们想要分析的实际计算:
int main(int argc, char *argv[])
{
// 初始化CUDA并获取设备/上下文
CUdevice device;
CUcontext context;
// 初始化CUDA驱动API
DRIVER_API_CALL(cuInit(0));
DRIVER_API_CALL(cuDeviceGet(&device, 0));
DRIVER_API_CALL(cuCtxCreate(&context, 0, device));
// 设置采样信息
SamplingInfo samplingInfo;
samplingInfo.device = device;
samplingInfo.context = context;
samplingInfo.terminate = 0;
// 默认为"inst_executed"或使用命令行参数
if (argc > 1) {
samplingInfo.eventName = argv[1];
} else {
samplingInfo.eventName = "inst_executed";
}
// 创建并启动采样线程
pthread_t sampling_thread;
pthread_create(&sampling_thread, NULL, sampling_func, &samplingInfo);
// 为向量加法分配内存
float *d_A, *d_B, *d_C;
float *h_A, *h_B, *h_C;
size_t size = VECTOR_SIZE * sizeof(float);
// 分配并初始化主机内存
h_A = (float *)malloc(size);
h_B = (float *)malloc(size);
h_C = (float *)malloc(size);
// 初始化向量
for (int i = 0; i < VECTOR_SIZE; i++) {
h_A[i] = i * 1.0f;
h_B[i] = i * 2.0f;
}
// 分配设备内存
RUNTIME_API_CALL(cudaMalloc((void **)&d_A, size));
RUNTIME_API_CALL(cudaMalloc((void **)&d_B, size));
RUNTIME_API_CALL(cudaMalloc((void **)&d_C, size));
// 将主机内存复制到设备
RUNTIME_API_CALL(cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice));
RUNTIME_API_CALL(cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice));
// 多次启动内核以便有时间采样
dim3 threadsPerBlock(256);
dim3 blocksPerGrid((VECTOR_SIZE + threadsPerBlock.x - 1) / threadsPerBlock.x);
for (int i = 0; i < ITERATIONS; i++) {
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, VECTOR_SIZE);
}
// 确保所有内核都完成
RUNTIME_API_CALL(cudaDeviceSynchronize());
// 通知采样线程终止
samplingInfo.terminate = 1;
// 等待采样线程完成
pthread_join(sampling_thread, NULL);
// 清理并退出
free(h_A);
free(h_B);
free(h_C);
RUNTIME_API_CALL(cudaFree(d_A));
RUNTIME_API_CALL(cudaFree(d_B));
RUNTIME_API_CALL(cudaFree(d_C));
DRIVER_API_CALL(cuCtxDestroy(context));
return 0;
}
此代码的关键方面:
- 长期运行的工作负载:内核运行2000次以确保我们有足够时间收集样本
- 线程协调:主线程在计算完成时通知采样线程
- 简单测试内核:使用向量加法作为采样的测试用例
3. 向量加法内核
__global__ void vecAdd(const float *A, const float *B, float *C, int numElements)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < numElements)
C[i] = A[i] + B[i];
}
这个简单的内核将两个向量相加。虽然不是计算密集型的,但重复运行为我们提供了一致的工作负载来监控。
运行教程
-
构建示例:
-
使用默认事件(指令计数)运行:
-
尝试不同的事件:
理解输出
使用默认的"inst_executed"事件运行时,您将看到如下输出:
inst_executed: 0
inst_executed: 25600000
inst_executed: 51200000
inst_executed: 76800000
inst_executed: 102400000
...
inst_executed: 4582400000
inst_executed: 4608000000
每行表示: 1. 被采样事件的名称 2. 采样时该事件的累积计数
在这种情况下,我们看到GPU执行的指令总数,随着我们的内核运行而稳步增加。规律的增量显示我们的工作负载在时间上执行得一致。
可采样的可用事件
不同的GPU支持不同的事件。您可能想要采样的一些常见事件包括:
inst_executed
:执行的指令branch
:执行的分支指令divergent_branch
:发散分支指令active_cycles
:至少一个warp活跃的周期