CUPTI 时间戳回调教程

完整的 GitHub 仓库和教程请访问 https://github.com/eunomia-bpf/cupti-tutorial。

简介

准确测量 CUDA 操作的执行时间对于性能优化至关重要。虽然基于 CPU 的计时可以给出近似结果，但 GPU 时间戳提供了操作在设备上实际执行时间的精确测量。本教程演示如何使用 CUPTI 回调来收集 CUDA 操作的 GPU 时间戳，为内存传输和内核执行提供准确的时序信息。

学习内容

如何为 CUDA 运行时 API 函数设置 CUPTI 回调
在操作开始和结束时收集精确的 GPU 时间戳
测量内存传输和内核的执行时间
分析完整 CUDA 工作流的性能

理解 GPU 时间戳

GPU 时间戳在几个重要方面与 CPU 时间戳不同：

它们从 GPU 的角度测量时间
它们提供纳秒级精度
它们准确捕获操作在设备上执行的时间，而不仅仅是提交的时间
它们对于理解 GPU 操作的真实性能特征至关重要

代码详解

1. 设置回调系统

首先，我们需要初始化 CUPTI 并注册我们的回调函数：

int main(int argc, char *argv[])
{
    // 初始化 CUDA
    RUNTIME_API_CALL(cudaSetDevice(0));

    // 订阅 CUPTI 回调
    CUpti_SubscriberHandle subscriber;
    CUPTI_CALL(cuptiSubscribe(&subscriber, (CUpti_CallbackFunc)getTimestampCallback, &traceData));

    // 启用 CUDA 运行时的回调
    CUPTI_CALL(cuptiEnableDomain(1, subscriber, CUPTI_CB_DOMAIN_RUNTIME_API));

    // 运行我们的测试工作负载
    runTest();

    // 取消订阅回调
    CUPTI_CALL(cuptiUnsubscribe(subscriber));

    return 0;
}

这设置了 CUPTI，在调用 CUDA 运行时 API 函数时调用我们的 getTimestampCallback 函数。

2. 回调函数

示例的核心是收集时间戳的回调函数：

void CUPTIAPI getTimestampCallback(void *userdata, CUpti_CallbackDomain domain,
                                  CUpti_CallbackId cbid, const CUpti_CallbackData *cbInfo)
{
    RuntimeApiTrace_t *traceData = (RuntimeApiTrace_t*)userdata;

    // 只处理运行时 API 回调
    if (domain != CUPTI_CB_DOMAIN_RUNTIME_API) return;

    // 我们对内存传输和内核启动感兴趣
    if ((cbid == CUPTI_RUNTIME_TRACE_CBID_cudaMemcpy_v3020) ||
        (cbid == CUPTI_RUNTIME_TRACE_CBID_cudaLaunch_v3020) ||
        (cbid == CUPTI_RUNTIME_TRACE_CBID_cudaDeviceSynchronize_v3020)) {

        if (cbInfo->callbackSite == CUPTI_API_ENTER) {
            // 在 API 调用开始时记录信息
            uint64_t timestamp;
            CUPTI_CALL(cuptiDeviceGetTimestamp(cbInfo->context, &timestamp));

            // 存储函数名
            traceData->functionName = cbInfo->functionName;
            traceData->startTimestamp = timestamp;

            // 对于内存传输，捕获大小和方向
            if (cbid == CUPTI_RUNTIME_TRACE_CBID_cudaMemcpy_v3020) {
                cudaMemcpy_v3020_params *params = 
                    (cudaMemcpy_v3020_params*)cbInfo->functionParams;
                traceData->memcpyBytes = params->count;
                traceData->memcpyKind = params->kind;
            }
        }
        else if (cbInfo->callbackSite == CUPTI_API_EXIT) {
            // 在 API 调用结束时记录信息
            uint64_t timestamp;
            CUPTI_CALL(cuptiDeviceGetTimestamp(cbInfo->context, &timestamp));

            // 计算持续时间
            traceData->gpuTime = timestamp - traceData->startTimestamp;

            // 打印时序信息
            printTimestampData(traceData);
        }
    }
}

此函数的关键方面： 1. 它过滤我们想要计时的特定 CUDA 函数 2. 在 API 进入时，它记录函数名、开始时间戳和任何相关参数 3. 在 API 退出时，它记录结束时间戳并计算持续时间 4. 对于内存传输，它捕获大小和方向

3. 测试工作负载

为了演示计时，我们运行一个简单的向量加法工作流：

void runTest()
{
    int N = 50000;
    size_t size = N * sizeof(float);
    float *h_A, *h_B, *h_C;
    float *d_A, *d_B, *d_C;

    // 分配主机内存
    h_A = (float*)malloc(size);
    h_B = (float*)malloc(size);
    h_C = (float*)malloc(size);

    // 初始化主机数组
    for (int i = 0; i < N; i++) {
        h_A[i] = rand() / (float)RAND_MAX;
        h_B[i] = rand() / (float)RAND_MAX;
    }

    // 分配设备内存
    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));

    // 启动内核（这将被计时）
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 将结果传输回主机（这将被计时）
    RUNTIME_API_CALL(cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost));

    // 同步以确保所有操作完成
    RUNTIME_API_CALL(cudaDeviceSynchronize());

    // 清理
    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));
}

此工作流包括： 1. 两个主机到设备的内存传输 2. 一个内核执行 3. 一个设备到主机的内存传输 4. 一个设备同步

4. 显示时序结果

结果以格式化表格显示：

void printTimestampData(RuntimeApiTrace_t *traceData)
{
    static bool headerPrinted = false;

    if (!headerPrinted) {
        printf("\n开始时间戳/GPU时间以纳秒报告\n\n");
        printf("%-15s %-24s %-15s %-8s %-10s\n", "名称", "开始时间", "GPU时间", "字节数", "类型");
        headerPrinted = true;
    }

    // 打印函数名和时序信息
    printf("%-15s %-24llu %-15llu", 
           traceData->functionName,
           (unsigned long long)traceData->startTimestamp,
           (unsigned long long)traceData->gpuTime);

    // 对于内存传输，打印大小和方向
    if (strcmp(traceData->functionName, "cudaMemcpy") == 0) {
        printf(" %-8lu %-10s", 
               (unsigned long)traceData->memcpyBytes,
               getMemcpyKindString((cudaMemcpyKind)traceData->memcpyKind));
    } else {
        printf(" %-8s %-10s", "不适用", "不适用");
    }

    printf("\n");
}

运行教程

构建示例：
```
make
```
运行时间戳收集器：
```
./callback_timestamp
```

理解输出

示例产生类似的输出：

开始时间戳/GPU时间以纳秒报告

名称            开始时间                GPU时间        字节数   类型
cudaMemcpy      123456789012            5432           200000   主机到设备
cudaMemcpy      123456794444            5432           200000   主机到设备
VecAdd          123456799876            10864          不适用    不适用
cudaMemcpy      123456810740            5432           200000   设备到主机
cudaDeviceSync  123456816172            0              不适用    不适用

让我们分析这个输出：

开始时间：操作开始时的 GPU 时间戳（以纳秒为单位）
GPU 时间：GPU 上操作的持续时间（以纳秒为单位）
字节数：对于内存传输，传输的数据量
类型：对于内存传输，方向（主机到设备或设备到主机）

从这个输出我们可以看到： - 内存传输每次大约需要 5.4 微秒 - 内核执行大约需要 10.8 微秒 - 整个工作流在大约 27 微秒内完成

性能洞察

这个时序数据揭示了 CUDA 性能的几个重要方面：

内存传输开销：内存传输可能是一个重要的瓶颈
内核执行时间：GPU 上的实际计算时间
操作序列：工作流中操作的顺序和时序

优化机会识别

基于时序数据，我们可以识别：

内存带宽利用率：

float bandwidth = (bytes / 1e9) / (time_seconds);
printf("内存带宽：%.1f GB/s\n", bandwidth);

计算密度：

float computeRatio = kernelTime / totalTime;
printf("计算时间比例：%.1f%%\n", computeRatio * 100);

API 开销：

float apiOverhead = (apiExitTime - apiEnterTime) - gpuTime;
printf("API 开销：%.2f 微秒\n", apiOverhead / 1000.0);

高级时序分析

多流分析

对于复杂的应用程序，您可能需要分析多个 CUDA 流：

struct StreamTimeline {
    int streamId;
    std::vector<OperationTiming> operations;

    void analyzeOverlap() {
        // 分析流之间的重叠
        for (size_t i = 0; i < operations.size() - 1; i++) {
            if (operations[i].endTime > operations[i+1].startTime) {
                printf("流 %d：操作重叠检测\n", streamId);
            }
        }
    }
};

内存传输模式

分析内存访问模式：

class MemoryAnalyzer {
private:
    std::map<void*, size_t> allocationSizes;
    std::vector<MemoryTransfer> transfers;

public:
    void analyzePattern() {
        size_t totalHostToDevice = 0;
        size_t totalDeviceToHost = 0;

        for (const auto& transfer : transfers) {
            if (transfer.kind == cudaMemcpyHostToDevice) {
                totalHostToDevice += transfer.bytes;
            } else if (transfer.kind == cudaMemcpyDeviceToHost) {
                totalDeviceToHost += transfer.bytes;
            }
        }

        printf("总主机到设备传输：%zu MB\n", totalHostToDevice / (1024*1024));
        printf("总设备到主机传输：%zu MB\n", totalDeviceToHost / (1024*1024));
        printf("数据传输比率：%.2f\n", 
               (double)totalDeviceToHost / totalHostToDevice);
    }
};

最佳实践

时序精度

使用 GPU 时间戳：始终使用 cuptiDeviceGetTimestamp 而不是 CPU 时钟
避免同步开销：最小化不必要的同步点
考虑时钟域：GPU 和 CPU 时钟可能不同步

性能影响

最小化回调开销：只跟踪关键操作
批量数据收集：避免在回调中进行昂贵的操作
后处理分析：将详细分析推迟到执行后

数据验证

bool validateTimingData(const TimingResult& result) {
    // 检查合理的时间范围
    if (result.duration < 0 || result.duration > MAX_REASONABLE_TIME) {
        return false;
    }

    // 验证内存传输速率
    if (result.isMemoryTransfer) {
        float bandwidth = calculateBandwidth(result);
        if (bandwidth > MAX_MEMORY_BANDWIDTH) {
            printf("警告：内存带宽超出预期：%.1f GB/s\n", bandwidth);
            return false;
        }
    }

    return true;
}

这个时间戳回调教程为精确测量和分析 CUDA 应用程序性能提供了强大的基础。

Share on Share on