[cuda graph 系列] cuda graph基本介绍

好的！我来简单解释一下 CUDA 图（CUDA Graph），它是 CUDA 编程中的一个优化工具，可以帮助提高 GPU 程序的效率，即使你对 CUDA 不熟悉也能理解。

什么是 CUDA Graph？

CUDA Graph 是 NVIDIA CUDA 编程模型中用来优化 GPU 程序性能的一种技术。它允许我们把一系列连续的 GPU 操作（例如计算和数据传输）连接成一个“图”结构，并一次性提交给 GPU 来执行。这样可以减少 CPU 和 GPU 之间的通信开销，让 GPU 更流畅、高效地工作。

为什么需要 CUDA Graph？

在传统的 GPU 编程中，每次要让 GPU 执行任务时，CPU 必须发送一个指令，并等待 GPU 完成任务后再发送下一个指令。这种“一个一个发送”的方式效率较低，因为 CPU 和 GPU 之间来回通信的开销会导致时间浪费。

CUDA Graph 通过把多个任务组成一个图结构，一次性提交给 GPU，这样 GPU 可以自己安排这些任务的执行顺序，而不需要每次都等待 CPU 的新指令，从而大大提高了效率。

CUDA Graph 的组成

CUDA Graph 主要由以下两部分组成：

1. 节点（Node）：图中的每个“节点”表示一个任务，比如一个计算任务或一个数据传输任务。
2. 边（Edge）：图中的每条“边”表示任务之间的依赖关系。例如，任务 A 必须在任务 B 之前完成，那么就会有一条从 A 到 B 的边。

通过这种图结构，GPU 能够按顺序执行任务，并且可以在不影响依赖关系的前提下并行执行多个任务。

CUDA Graph 如何工作？

CUDA Graph 的基本流程如下：

1. 创建图：我们先定义好需要 GPU 执行的所有任务，并确定它们之间的顺序和依赖关系。
2. 录制图（Recording）：把这些任务记录成一个图结构。
3. 执行图（Execution）：把这个图提交给 GPU，让 GPU 按照图中的任务顺序来执行。

优点

使用 CUDA Graph 有以下几个优点：

• 减少 CPU 和 GPU 的通信开销：减少了 CPU 向 GPU 发送指令的次数，使 GPU 能够连续执行任务。
• 提高并行性：GPU 可以通过并行处理多个没有依赖关系的任务，进一步提高效率。
• 优化长时间运行的任务：对于那些需要重复执行的任务，可以把它们录制成图并反复运行，从而省去每次都要重新提交的开销。

示例

假设我们有三个任务，A -> B -> C，其中任务 A 和 B 之间有依赖关系，B 和 C 之间也有依赖关系。那么我们可以用 CUDA Graph 来把这三个任务串起来，交给 GPU 一次性执行。

这种方式会比每个任务单独提交给 GPU 的方式更快，特别是在任务量大、任务依赖多的情况下。

总结

CUDA Graph 是一种加速 GPU 程序执行的技术，它通过将任务组织成图结构，减少了 CPU 和 GPU 之间的通信开销，并且允许 GPU 自己安排任务的执行顺序，从而显著提高了性能。

CUDA Graph 的基本使用流程

使用 CUDA Graph 的主要步骤包括创建图、录制图、执行图等。

CUDA Graph 基本使用步骤

在 CUDA 中使用 CUDA Graph 大致可以分为以下几步：

1. 创建和开始录制图：定义一个 CUDA Graph 对象，准备好图的录制。
2. 将操作加入到图中：在图的录制过程中，将各种 GPU 操作（比如内存拷贝、核函数调用等）加入到图中。
3. 结束图的录制：结束录制并生成一个可执行的图对象。
4. 执行图：将图提交给 GPU 执行。
5. 清理资源：释放图和相关资源。

代码示例

我们来看一个简单的代码示例，展示如何在 CUDA 中使用 CUDA Graph。这里的例子假设你已经对 CUDA 编程有一些基本了解，例如如何编写核函数。

示例代码：使用 CUDA Graph 来加速简单的 GPU 操作

假设我们要做一个向量加法 C = A + B，这类任务在传统 CUDA 中需要分别调用内存分配、内存拷贝、核函数执行等步骤。使用 CUDA Graph，我们可以将这些操作录制成一个图，并一次性提交给 GPU。

#include <cuda_runtime.h>
#include <iostream>

// 核函数定义：向量加法
__global__ void vectorAdd(const float* A, const float* B, float* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

int main() {
    int N = 1 << 20;  // 向量长度
    size_t size = N * sizeof(float);

    // 分配并初始化主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);
    for (int i = 0; i < N; i++) {
        h_A[i] = 1.0f;
        h_B[i] = 2.0f;
    }

    // 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    // 定义 CUDA Graph 对象和图实例
    cudaGraph_t graph;
    cudaGraphExec_t instance;

    // 开始录制 CUDA Graph
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

    // 将操作加入到图中
    cudaMemcpyAsync(d_A, h_A, size, cudaMemcpyHostToDevice, stream);
    cudaMemcpyAsync(d_B, h_B, size, cudaMemcpyHostToDevice, stream);

    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    vectorAdd<<<numBlocks, blockSize, 0, stream>>>(d_A, d_B, d_C, N);

    cudaMemcpyAsync(h_C, d_C, size, cudaMemcpyDeviceToHost, stream);

    // 结束录制并生成可执行图
    cudaStreamEndCapture(stream, &graph);
    cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);

    // 执行图
    cudaGraphLaunch(instance, stream);
    cudaStreamSynchronize(stream);

    // 清理资源
    cudaGraphDestroy(graph);
    cudaGraphExecDestroy(instance);
    cudaStreamDestroy(stream);

    // 验证结果
    bool success = true;
    for (int i = 0; i < N; i++) {
        if (h_C[i] != 3.0f) {
            success = false;
            break;
        }
    }
    std::cout << (success ? "Success" : "Failure") << std::endl;

    // 释放内存
    free(h_A);
    free(h_B);
    free(h_C);
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    return 0;
}

代码解释

1. 数据准备：我们在主机上分配和初始化输入向量 h_A 和 h_B，以及输出向量 h_C。
2. 设备内存分配：在 GPU 上为 d_A、d_B、d_C 分配内存，用于存储输入和输出数据。
3. CUDA Graph 创建和录制：
   • 通过 cudaStreamBeginCapture 开始在 stream 流中捕获操作，准备录制图。
   • 将内存拷贝和核函数调用操作加入到图中。
4. 结束录制：
   • 通过 cudaStreamEndCapture 结束图的录制，生成 graph 对象。
   • 使用 cudaGraphInstantiate 创建图的可执行实例 instance。
5. 图的执行：使用 cudaGraphLaunch 将图提交给 GPU 执行。
6. 同步和结果检查：等待图执行完成，验证结果是否正确。
7. 清理资源：释放所有分配的资源，包括图和设备内存等。

关键函数

• cudaStreamBeginCapture 和 cudaStreamEndCapture：用于开始和结束图的录制。
• cudaGraphInstantiate：将录制的图转换为可执行实例。
• cudaGraphLaunch：在 GPU 上启动图的执行。
• cudaGraphDestroy 和 cudaGraphExecDestroy：用于清理图对象和执行实例。

使用 CUDA Graph 的好处

1. 减少开销：将一系列操作打包成一个图，减少了 CPU 和 GPU 的同步等待。
2. 性能提升：适合反复执行的任务，可以更好地利用 GPU 并行性和流水线执行能力。
3. 代码简化：让复杂的任务序列更易于管理，特别是有大量依赖关系的复杂任务。

总结

通过 CUDA Graph，可以将多个 GPU 操作封装成一个图，提高了执行效率和性能，尤其适合那些需要反复执行的计算密集型任务。