CUDA

CUDA

CUDA Event精确测量GPU kernel的耗时

CUDA Event 是 NVIDIA 提供的 GPU 时间戳工具，直接记录 GPU 硬件的时间，解决了异步问题：

  1. 同步机制：
    cudaEventRecord() 会在 GPU 的特定时间点插入事件标记，cudaEventSynchronize() 会等待 GPU 执行到该事件。
  2. 精确测量：
    计算的是 GPU 实际执行核函数的耗时（而非 CPU 时间）。

示例代码（正确方法）：

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start);            // 记录 GPU 开始时间
vectorAdd<<<grid, block>>>(a, b, c, N);
cudaEventRecord(stop);             // 记录 GPU 结束时间
cudaEventSynchronize(stop);        // 等待 GPU 执行完成

float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);  // 计算 GPU 时间差
printf("Kernel time: %.3f ms\n", milliseconds);

cudaEventDestroy(start);
cudaEventDestroy(stop);

替代方案（不推荐）

如果坚持不用 CUDA Event，可以 强制同步 GPU，但仍有缺陷：

vectorAdd<<<grid, block>>>(a, b, c, N);
cudaDeviceSynchronize();  // 阻塞 CPU，等待 GPU 所有任务完成
  • 缺点：
    • 会测量 所有未完成的 GPU 操作（包括其他核函数、内存拷贝等），而不仅是目标核函数。
    • 不如 CUDA Event 精准和灵活。
方法 适用场景 精度 缺点
CUDA Event 精确测量 GPU 核函数/内存操作 高（微秒级） 需额外代码
cudaDeviceSynchronize() 粗略测量 GPU 整体耗时 低 无法隔离单个核函数
clock() 仅测量 CPU 时间 不适用 GPU 完全错误

结论：

• 必须使用 CUDA Event 测量核函数时间（如 vectorAdd）。
• Nsight 工具中的 cudaLaunchKernel 时间也是通过 GPU 事件机制捕获的，与手动使用 cudaEvent 结果一致。

CUDA Cores, Blocks, Threads

在 CUDA 架构中，一个 CUDA Core 并不直接对应一个 Block，而是 Block 中的线程（Threads）会被分配到 CUDA Core 上执行。以下是详细解释：

1. CUDA Core 与 Block 的关系
• CUDA Core：

• 是 GPU 最基本的计算单元，负责执行单个线程的算术逻辑运算（如浮点加法、乘法）。

• 例如：你的 RTX 3070 Ti 有 5888 个 CUDA Cores（46 SM × 128 Cores/SM）。

• Block：

• 是逻辑上的线程组（如 <<<grid, block>>> 中的 block），包含多个线程（最多 1024 个）。

• Block 会被分配到 SM（Streaming Multiprocessor）上执行，而非直接映射到 CUDA Core。

2. 执行机制

  1. Block 分配到 SM：
    • 每个 SM 可以同时驻留多个 Block（如 RTX 3070 Ti 的每个 SM 支持 最多 16 个 Block）。

    • 你的 GPU 有 46 个 SM，因此理论上可同时运行 46 × 16 = 736 个 Block（实际受资源限制）。

  2. 线程分配到 CUDA Core：
    • 每个 SM 中的线程以 Warp（32 线程） 为单位调度到 CUDA Core 上执行。

    • 例如：一个 Block 有 256 线程 → 会被分成 8 个 Warp（256/32）。

  3. 动态调度：
    • SM 的 Warp 调度器会动态将就绪的 Warp 分配到空闲的 CUDA Core 上，隐藏内存延迟。

    • 一个 CUDA Core 可以先后执行多个线程（非一对一绑定）。

3. 为什么不是一对一绑定？
• 并行性需求：

• GPU 通过 大量线程切换 来隐藏内存访问延迟（如全局内存访问需要数百周期）。

• 若一个 CUDA Core 只处理一个线程，计算资源会因等待内存而闲置。

• 硬件设计：

• 每个 SM 的 CUDA Core 数量（如 128）远小于其支持的线程数（如 1536），因此必须复用。

4. 直观类比

概念 类比 说明
CUDA Core 工厂工人 实际执行任务的工人（计算单元）。
Block 生产小组 一组工人（线程）共同完成一个任务。
SM 车间 可同时管理多个生产小组（Block）。
Warp 流水线批次 工人每次处理 32 个任务（线程），轮流切换。

5. 如何验证？
通过 deviceQuery 输出中的以下参数计算：
• 每个 SM 的 CUDA Cores：128（RTX 3070 Ti）。

• 每个 SM 的最大线程数：1536。

• 线程复用比例：1536 / 128 = 12 线程/CUDA Core（通过时间片轮转实现）。

6. 优化建议

  1. Block 大小选择：
    • 每个 Block 至少 128~256 线程，以填充 SM 的 Warp 调度容量。

    • 示例：

    dim3 block(256);  // 每个 Block 256 线程
    dim3 grid((N + block.x - 1) / block.x);  // 覆盖所有数据
    kernel<<<grid, block>>>(...);

  2. 资源限制：
    • 每个 Block 的共享内存和寄存器使用会影响 SM 中驻留的 Block 数量。

  3. Occupancy（占用率）：
    • 使用 CUDA Occupancy Calculator 工具，最大化 SM 的线程利用率。

总结
• 一个 CUDA Core 不固定对应一个 Block，而是通过 Warp 调度动态执行多个线程。
• 关键目标：让每个 SM 的 CUDA Core 始终有可执行的线程（通过足够多的 Block 和线程）。
• 硬件参数：以 RTX 3070 Ti 为例，46 个 SM × 1536 线程/SM = 70,656 个并发线程潜力。