CUDA Graphs学习与实验

CUDA Graphs学习与实验

引用

CSDN

1.

https://blog.csdn.net/m0_61864577/article/details/142846235

CUDA图（CUDA Graphs）为CUDA引入了一种全新的工作提交模型。它允许将一系列操作（如内核启动）以图的形式表示，并通过依赖关系将这些操作连接起来。这种图的定义过程与其执行过程是分开的，这意味着我们可以提前定义好一个图，然后多次重复执行它。

这种定义与执行的分离带来了多方面的优化：

1. 降低CPU启动开销：相比传统的流（streams）方式，由于大量的设置工作已经在图的定义和实例化阶段完成，实际执行时的CPU开销明显减少。
2. 全局优化机会：通过将整个工作流程以图的形式呈现给CUDA，CUDA有机会对整个流程进行优化。这在逐步提交工作的流机制中是无法实现的，因为流机制只能看到局部的、片段式的工作提交。

流机制中的问题

在传统的流中，当你向流中放置一个内核时，主机驱动程序需要执行一系列操作来准备在GPU上执行该内核。这些操作包括设置内核参数、配置执行环境等。对于执行时间较短的GPU内核，这些准备工作的开销可能占到总执行时间的很大一部分，从而降低了整体效率。

CUDA图的工作提交分为三个阶段

1. 定义（Definition）：
   在这个阶段，程序创建一个包含操作及其依赖关系的图。开发者描述需要执行的操作（如内核函数）以及这些操作之间的先后顺序或并行关系。

2. 实例化（Instantiation）：
   在定义完成后，CUDA对图进行实例化。实例化过程包括：

• 快照：对图模板进行捕获，生成一个具体的可执行图结构。
• 验证：检查图的正确性，确保所有的操作和依赖关系都是有效的。
• 预处理：执行大部分的设置和初始化工作，目的是尽可能减少在实际执行时需要完成的工作量。
  实例化的结果是一个可执行图（executable graph）。

3. 执行（Execution）：
   已实例化的可执行图可以像普通的CUDA工作一样被提交到流中执行。重要的是，这个可执行图可以被多次执行，而无需每次都重新实例化。这大大提高了执行的效率，特别是在需要重复执行相同操作的情况下。

CUDA图的优势

• 性能提升：通过减少CPU的启动开销，特别是在需要频繁启动小型内核的情况下，CUDA图能够显著提升性能。
• 优化执行：由于CUDA能够提前知道整个工作流程，它可以进行全局优化。例如，它可以重新排列操作以提高并行性，或者优化内存传输以减少延迟。
• 简化编程模型：开发者可以以更直观的方式描述计算任务，而无需手动管理复杂的依赖关系和同步机制。

举例说明

假设我们有一系列需要按特定顺序执行的内核操作。在传统的流机制中，我们需要：

• 为每个内核启动，都要进行一次完整的设置和启动过程。
• 手动管理这些内核之间的依赖关系，确保它们按正确的顺序执行。

使用CUDA图后，我们可以：

• 一次性地定义所有的内核操作和它们的依赖关系。
• 实例化后，CUDA会处理好所有的设置和依赖关系。
• 执行时，只需简单地启动可执行图即可。

结论

CUDA图为GPU计算提供了更高效、更灵活的工作提交方式。通过预先定义和实例化计算图，CUDA能够减少CPU的开销，并利用全局信息对执行进行优化。这对于需要高性能计算的应用，尤其是包含大量小型、短时内核的应用，具有重要意义。

测试方案

测试代码

#include <iostream>
#include <cuda_runtime.h>
#include <stdio.h>
#include <assert.h>
#include <cuda.h>

#define CHECK_CUDA(call)                                            \
    do {                                                            \
        cudaError_t err = call;                                    \
        if (err != cudaSuccess) {                                 \
            std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \
            std::cerr << " code=" << err << " (" << cudaGetErrorString(err) << ")" << std::endl; \
            exit(EXIT_FAILURE);                                    \
        }                                                         \
    } while (0)

__global__ void Kernel1(float *a, float *b, float *c, float *d) {
    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
    a[tid] = 1; b[tid] = 2; c[tid] = 3; d[tid] = 0;
    if (tid == 0) {
        printf("Kernel1\n");
    }
}

__global__ void Kernel2(float *a, float *b, float *c, float *d) {
    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
    a[tid] += 1;
    if (tid == 0) {
        printf("Kernel2\n");
    }
}

__global__ void Kernel3(float *a, float *b, float *c, float *d) {
    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
    b[tid] += 2;
    if (tid == 0) {
        printf("Kernel3\n");
    }
}

__global__ void Kernel4(float *a, float *b, float *c, float *d) {
    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
    c[tid] += 3;
    if (tid == 0) {
        printf("Kernel4\n");
    }
}

__global__ void Kernel5(float *a, float *b, float *c, float *d) {
    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
    d[tid] = a[tid] + b[tid] + c[tid];
    if (tid == 0) {
        printf("Kernel5\n");
    }
}

void CUDART_CB HostFunc(void *data) {
    float *d = static_cast<float*>(data);
    for (int i = 0; i < thread_size; i++) {
        d[i] += 1;
    }
    printf("HostFunc\n");
}

int run(bool graph_mode) {
    int deviceid = 0;
    int block_count = 1;
    int block_size = 8;
    int thread_size = block_count * block_size;
    int total_count = thread_size * sizeof(float);

    cudaStream_t stream[3];
    cudaEvent_t event[3];

    CHECK_CUDA(cudaSetDevice(deviceid));
    for (int i = 0; i < 3; i++) {
        CHECK_CUDA(cudaStreamCreate(&stream[i]));
        CHECK_CUDA(cudaEventCreate(&event[i]));
    }

    float *a, *b, *c, *d;
    CHECK_CUDA(cudaMallocManaged(&a, total_count));
    CHECK_CUDA(cudaMallocManaged(&b, total_count));
    CHECK_CUDA(cudaMallocManaged(&c, total_count));
    CHECK_CUDA(cudaMallocManaged(&d, total_count));

    cudaGraph_t graph;
    if (graph_mode) {
        CHECK_CUDA_DRV_API(cuGraphCreate(&graph, 0));
        CHECK_CUDA(cudaStreamBeginCapture(stream[0], cudaStreamCaptureModeGlobal));
    }

    Kernel1<<<block_count, block_size, 0, stream[0]>>>(a, b, c, d);
    CHECK_CUDA(cudaEventRecord(event[0], stream[0]));

    CHECK_CUDA(cudaStreamWaitEvent(stream[1], event[0]));
    CHECK_CUDA(cudaStreamWaitEvent(stream[2], event[0]));

    Kernel2<<<block_count, block_size, 0, stream[0]>>>(a, b, c, d);
    Kernel3<<<block_count, block_size, 0, stream[1]>>>(a, b, c, d);
    CHECK_CUDA(cudaEventRecord(event[1], stream[1]));

    Kernel4<<<block_count, block_size, 0, stream[2]>>>(a, b, c, d);
    CHECK_CUDA(cudaEventRecord(event[2], stream[2]));

    CHECK_CUDA(cudaStreamWaitEvent(stream[0], event[1]));
    CHECK_CUDA(cudaStreamWaitEvent(stream[0], event[2]));

    Kernel5<<<block_count, block_size, 0, stream[0]>>>(a, b, c, d);
    CHECK_CUDA(cudaLaunchHostFunc(stream[0], HostFunc, (void*)d));

    if (graph_mode) {
        CHECK_CUDA(cudaStreamEndCapture(stream[0], &graph));
        cudaGraphExec_t graphExec;
        CHECK_CUDA(cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
        CHECK_CUDA(cudaGraphLaunch(graphExec, 0));
        CHECK_CUDA(cudaDeviceSynchronize());
        CHECK_CUDA_DRV_API(cuGraphDebugDotPrint(graph, "graph.dot", 0));
        CHECK_CUDA(cudaGraphExecDestroy(graphExec));
        CHECK_CUDA(cudaGraphDestroy(graph));
    } else {
        CHECK_CUDA(cudaStreamSynchronize(stream[0]));
    }

    for (int i = 0; i < thread_size; i++) {
        printf("%6.2f\n", d[i]);
    }

    CHECK_CUDA(cudaFree(a));
    CHECK_CUDA(cudaFree(b));
    CHECK_CUDA(cudaFree(c));
    CHECK_CUDA(cudaFree(d));
    return 0;
}

int main(int argc, char *argv[]) {
    int mode = atoi(argv[1]);
    if (mode == 0) {
        printf("normal mode\n");
        run(0);
    } else {
        printf("graph mode\n");
        run(1);
    }
}

输出

normal mode
Kernel1
Kernel2
Kernel3
Kernel4
Kernel5
HostFunc
 13.00
 13.00
 13.00
 13.00
 13.00
 13.00
 13.00
 13.00
graph mode
Kernel1
Kernel2
Kernel3
Kernel4
Kernel5
HostFunc
 13.00
 13.00
 13.00
 13.00
 13.00
 13.00
 13.00
 13.00