CUPTI（CUDA Profiling Tools Interface）是 NVIDIA 提供的底层 profiling 接口。Nsight Systems、Nsight Compute 等工具底层都会使用 CUPTI 获取 CUDA 程序的执行信息。

对于 tracing 来说，最核心的两套机制是：

• Activity API：异步记录执行过程中发生了什么。
• Callback API：在某个 CUDA API 正在执行时同步插入 callback，在某个事件触发前/后进入这个callback function

---

1. Activity API

Activity API 用于收集执行 trace，例如：

cudaLaunchKernel
cudaMemcpyAsync
kernel execution
memcpy execution

它的整体流程只有四步：

1. RegisterCallbacks
   ↓
告诉 CUPTI：
“buffer 怎么向我要、写好之后怎么交给我”

2. ActivityEnable
   ↓
告诉 CUPTI：
“我要记录哪些类型的事件”

3. 程序正常运行
   ↓
CUPTI 收集这些事件并生成 activity records

4. CUPTI 把 records 放入 buffer
   ↓
交给 profiler 解析

1.1 cuptiActivityRegisterCallbacks

首先需要注册两个 buffer callback：

cuptiActivityRegisterCallbacks(
    bufferRequested,
    bufferCompleted
);

这个函数本身不决定记录什么事件。

它只是告诉 CUPTI：

需要新的 buffer
    → 调用 bufferRequested()

一块 buffer 可以交给 profiler
    → 调用 bufferCompleted()

因此两个 callback 分工非常明确。

bufferRequested

CUPTI 需要一块新的空 buffer 时调用：

CUPTI
  │
  │ “给我一块空 buffer”
  ▼
bufferRequested()
  │
  ▼
返回一块 host memory

例如：

void bufferRequested(
    uint8_t **buffer,
    size_t *size,
    size_t *maxNumRecords)
{
    *size = 1024 * 1024;
    *buffer = (uint8_t *)malloc(*size); //这里用malloc,所以我们completed的时候不用size就可以free,因为malloc维护了metadata
    *maxNumRecords = 0;
}

这块 buffer 一般位于 CPU host memory。

之后 CUPTI 会把 activity records 填入其中。

---

bufferCompleted

当 CUPTI 决定把这块 buffer 交给 profiler 处理时，就调用：

CUPTI
  │
  │ filled buffer
  ▼
bufferCompleted()

例如 buffer 中可能已经包含：

┌──────────────────────────┐
│ Runtime API record       │
│ Kernel record            │
│ Runtime API record       │
│ Memcpy record            │
│ ...                      │
└──────────────────────────┘

在 bufferCompleted() 中，可以使用：

cuptiActivityGetNextRecord(...)

逐条读取：

record 1
record 2
record 3
...

所以两个 callback 的本质是一个 producer-consumer 协议：

                 empty buffer
Profiler ─────────────────────► CUPTI

                filled buffer
Profiler ◄───────────────────── CUPTI

bufferRequested 解决：

CUPTI 去哪里写数据？

bufferCompleted 解决：

CUPTI 写好的数据怎么交还给 profiler？

void CUPTIAPI bufferCompleted(
    CUcontext ctx,
    uint32_t streamId,
    uint8_t *buffer,
    size_t size,
    size_t validSize)
{
    // 1. 解析 CUPTI 写进去的 Activity records
    CUpti_Activity *record = nullptr;

    while (cuptiActivityGetNextRecord(
               buffer,
               validSize,
               &record) == CUPTI_SUCCESS) {

        // 2. 根据 record 类型处理
        switch (record->kind) {
        case CUPTI_ACTIVITY_KIND_KERNEL:
            // 处理 kernel 信息
            break;

        case CUPTI_ACTIVITY_KIND_MEMCPY:
            // 处理 memcpy 信息
            break;

        default:
            break;
        }
    }

    // 3. 检查是否有 record 因为 buffer 太小被丢掉
    size_t dropped = 0;
    cuptiActivityGetNumDroppedRecords(
        ctx, streamId, &dropped);

    // 4. 释放这块 buffer
    free(buffer);
}

---

1.2 cuptiActivityEnable

注册完 buffer 管理方式后，还需要告诉 CUPTI：

到底要记录什么？

这就是：

cuptiActivityEnable(kind);

例如：

cuptiActivityEnable(
    CUPTI_ACTIVITY_KIND_RUNTIME
);

表示记录 CUDA Runtime API，例如：

cudaMalloc
cudaMemcpyAsync
cudaLaunchKernel
...

如果再打开：

cuptiActivityEnable(
    CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL
);

就会记录 GPU kernel execution，例如：

kernel name
start timestamp
end timestamp
stream
device
correlation ID
...

因此：

cuptiActivityRegisterCallbacks()
    = 设置数据怎么交付

cuptiActivityEnable()
    = 设置收集哪些数据

这是 Activity API 最重要的两个概念。

---

1.3 Activity API 的完整流程

例如：

cuptiActivityRegisterCallbacks(
    bufferRequested,
    bufferCompleted
);

cuptiActivityEnable(
    CUPTI_ACTIVITY_KIND_RUNTIME
);

cuptiActivityEnable(
    CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL
);

kernel<<<...>>>();

cuptiActivityFlushAll(0);

整个过程是：

cuptiActivityRegisterCallbacks(...)
        │
        ▼
告诉 CUPTI：
“以后需要 buffer / buffer 写完时，调用我这两个函数”
        │
        ▼
cuptiActivityEnable(KERNEL)
cuptiActivityEnable(RUNTIME)
        │
        ▼
告诉 CUPTI：
“我要收集 Kernel、Runtime 这些 Activity”
        │
        ▼
CUDA 程序开始运行
        │
        ▼
CUPTI 需要地方保存 Activity 数据
        │
        ▼
bufferRequested()
        │
        ▼
你 malloc 一块 CPU buffer 给 CUPTI
        │
        ▼
kernel / runtime API 不断发生
        │
        ▼
CUPTI 把对应 Activity record 写进 buffer
        │
        ▼
buffer 满了 / CUPTI 决定交付 / 你调用 flush
        │
        ▼
bufferCompleted(buffer, validSize, ...)
        │
        ▼
这个 buffer 的所有权交回给你
        │
        ▼
cuptiActivityGetNextRecord(...)
        │
        ▼
逐个读出 Kernel / Runtime / Memcpy record
        │
        ▼
Profiler 保存 / 分析 trace
        │
        ▼
free(buffer)

最后的：

cuptiActivityFlushAll(0);

用于把还没有自动交付的 activity buffer 强制 flush 出来。只有满了的buffer才会自动提交，如果我profile结束还有一块未满的buffer我就要flush出来。flush会自动调用bufferCompleted。

因此可以把 Activity API 压缩成一句话：

先注册 buffer 的供应和回收方式，再指定需要记录的 activity 类型；程序运行时 CUPTI 异步收集这些事件，并通过 buffer 批量交给 profiler。

---

2. Callback API

Callback API 的设计和 Activity 完全不同。

Activity 是：

事件发生
    ↓
记录下来
    ↓
放进 buffer
    ↓
之后处理

Callback 是：

事件正在发生
    ↓
立刻调用你的函数

因此它是一种同步 interception 机制。

---

2.1 cuptiSubscribe

第一步是注册一个 callback handler：

CUpti_SubscriberHandle subscriber;

cuptiSubscribe(
    &subscriber, //获取subscriber
    ProfilerCallbackHandler, //我的回调函数
    userdata //调用时把一些CPU侧的自己的数据传进去
);

这里有两个不同的概念：

ProfilerCallbackHandler
    = 真正被调用的函数

subscriber
    = 这次 subscription 的 handle

可以理解成：

cuptiSubscribe
      │
      ├── callback = ProfilerCallbackHandler
      ├── userdata = ...
      │
      ▼
创建一个 subscription
      │
      ▼
返回 subscriber handle

---

2.2 cuptiEnableCallback

有了 subscriber 以后，还没有告诉 CUPTI：

到底监听哪个事件？

因此还需要：

cuptiEnableCallback(
    1, //开启callback
    subscriber,
    CUPTI_CB_DOMAIN_DRIVER_API, //打开的是Driver API（同一个cbid可能在多个API里都有，runtime是高层C++层API,Driver API是底层）
    CUPTI_DRIVER_TRACE_CBID_cuLaunchKernel  //监听cuLaunchKernel （对应cbid,是一个enum）
);

这里表达的是：

对 subscriber 这个 subscription
启用：
Domain = DRIVER_API
Callback ID = cuLaunchKernel

也就是：

当程序调用 cuLaunchKernel 时，通知这个 subscriber。

因为 subscriber 已经和：

ProfilerCallbackHandler

绑定，所以真正发生事件时：

cuLaunchKernel()
        │
        ▼
CUPTI 检查订阅
        │
        ▼
subscriber 对这个事件开启
        │
        ▼
ProfilerCallbackHandler(...)

因此：

cuptiSubscribe
    = 注册“事件发生时调用谁”

cuptiEnableCallback
    = 指定“监听哪些事件”

---

3. Domain 和 Callback ID

Callback API 中，每个事件由：

Domain + Callback ID

共同确定。

例如：

DRIVER_API
├── cuLaunchKernel
├── cuMemAlloc
└── ...

RUNTIME_API
├── cudaLaunchKernel
├── cudaMalloc
├── cudaMemcpy
└── ...

所以：

CUPTI_CB_DOMAIN_DRIVER_API

表示：

CUDA Driver API 这一类事件。

而：

CUPTI_DRIVER_TRACE_CBID_cuLaunchKernel

表示：

其中的 cuLaunchKernel。

因此 handler 的参数里面也会带：

CUpti_CallbackDomain domain,
CUpti_CallbackId callbackId

让同一个 handlder 可以处理多个不同事件。

---

4. Callback 的 ENTER 和 EXIT

Callback API 一个非常重要的特点是：

对一次 CUDA API 调用，handler 通常会在 API 入口和出口都被调用。

例如：

cuLaunchKernel(...)
      │
      ├── ENTER
      │     ↓
      │   Handler()
      │
      │
      │   真正执行 cuLaunchKernel
      │
      │
      └── EXIT
            ↓
          Handler()

两次调用的都是同一个 handler。

通过：

callbackData->callbackSite

判断是哪一次，并在ProfilerCallbackHandle里面过滤。

例如：

//在void CUPTIAPI ProfilerCallbackHandle 里面
if (callbackData->callbackSite == CUPTI_API_ENTER) {
    // API 执行之前
}

if (callbackData->callbackSite == CUPTI_API_EXIT) {
    // API 即将返回
}

两个值分别是：

CUPTI_API_ENTER
CUPTI_API_EXIT

---

5. ENTER / EXIT 不等于 GPU kernel 开始 / 结束

这是非常容易混淆的一点。

对于：

kernel<<<...>>>();

CPU 最终执行一个 kernel launch API。

Callback 观察的是：

CPU API 调用

而不是：

GPU kernel execution

例如：

CPU:

cuLaunchKernel
│ ENTER
│
│ 提交 kernel
│
│ EXIT
└──────────────────────────────► time


GPU:

             │ kernel starts
             │
             │===================│
                                 kernel ends

因为 kernel launch 通常是异步的，所以：

cuLaunchKernel EXIT

只表示：

CPU 的 launch API 返回了。

完全不意味着：

GPU kernel 已经执行结束。

GPU kernel 真正的 start/end timestamp 更适合通过 Activity API 获得。

---

6. Callback API 的完整流程

最小代码：

CUpti_SubscriberHandle subscriber;

cuptiSubscribe(
    &subscriber,
    ProfilerCallbackHandler,
    userdata
);

cuptiEnableCallback(
    1,
    subscriber,
    CUPTI_CB_DOMAIN_DRIVER_API,
    CUPTI_DRIVER_TRACE_CBID_cuLaunchKernel
);

kernel<<<...>>>();

执行过程：

cuptiSubscribe
      │
      ▼
注册 Handler
      │
      ▼
获得 subscriber
      │
      ▼
cuptiEnableCallback
      │
      ▼
指定监听 cuLaunchKernel
      │
      ▼
程序调用 cuLaunchKernel
      │
      ├── Handler(ENTER)
      │
      ├── 真正执行 API
      │
      └── Handler(EXIT)

所以 Callback API 可以概括成：

先注册一个 Handler 并得到 subscriber，再指定这个 subscriber 要监听哪些 CUDA 事件；事件真正发生时，CUPTI 会同步调用 Handler。

---

7. Activity 和 Callback 的区别

最终可以用这一张图理解 CUPTI tracing：

                       CUPTI Tracing
                             │
             ┌───────────────┴───────────────┐
             │                               │
         Activity                         Callback
             │                               │
       异步记录事件                       同步拦截事件
             │                               │
   RegisterCallbacks                     Subscribe
             │                               │
     管理 trace buffer                  注册 Handler
             │                               │
     ActivityEnable                    EnableCallback
             │                               │
    指定记录什么事件                   指定监听什么事件
             │                               │
       CUDA 程序运行                      CUDA API 发生
             │                               │
       产生 records                   Handler(ENTER)
             │                               │
    写入 activity buffer                 API 执行
             │                               │
     BufferCompleted                  Handler(EXIT)
             │
     GetNextRecord

二者最重要的区别是：