CUDA Graphs学习与实验

于 2024-10-11 10:03:58 发布
阅读量1.3k 收藏 8
点赞数 30
分类专栏： 代码片段 文章标签： GPGPU GPU CUDA 学习
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/m0_61864577/article/details/142846235
版权

CUDA Graphs学习与实验

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的开销，并利用全局信息对执行进行优化。这对于需要高性能计算的应用，尤其是包含大量小型、短时内核的应用，具有重要意义

一.参考链接

二.测试方案

请添加图片描述

三.测试代码

tee cuda_graph.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <iostream>
#include <chrono>
#include <thread>

#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)


#define CHECK_CUDA_DRV_API(call)                      \
  do {                              \
    CUresult err = call;                  \
    if (err != CUDA_SUCCESS) {                 \
        char *error_str=new char[1024];  \
        cuGetErrorString(err,(const char**)&error_str); \
        printf("[%s:%d] %s Error :%s!\n",__FILE__,__LINE__,#call,error_str); \
    }                                      \
  } 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");
    }
}

struct HostFuncParam
{
    float *a;
    float *b;
    float *c;
    float *d;
    int thread_size;
};

void CUDART_CB HostFunc(void *data){
    HostFuncParam *pstParam=(HostFuncParam*)data;
    for(int i=0;i<pstParam->thread_size;i++)
    {
        pstParam->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);
    HostFuncParam stParam;
    stParam.d=d;
    stParam.thread_size=thread_size;
    CHECK_CUDA(cudaLaunchHostFunc(stream[0], HostFunc, (void*)&stParam));
    
    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);
    }
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -lineinfo \
        -o cuda_graph cuda_graph.cu \
        -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
./cuda_graph 0
./cuda_graph 1
dot -Tpng graph.dot -o graph.png
  • 输出
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

请添加图片描述

CUDA：Cuda Graphs图实例
希望我的博客，能帮上你解决学习中工作中所遇到的问题
12-21 381
CUDA：Cuda Graphs图实例
pytorch与深度学习
数字人生
04-17 1881
PyTorch是一个由Facebook AI Research Team开发的开源深度学习库，它提供了一个灵活的环境和丰富的API，用于快速且方便地构建、训练和部署深度学习模型。综上所述，PyTorch 凭借其灵活的动态图机制、强大的张量运算与自动求导、丰富的深度学习模块以及活跃的社区支持，已成为深度学习研究与应用开发的首选工具之一。2. 优秀的社区和生态系统：有着活跃的社区和广泛的支持，PyTorch提供了大量的预训练模型和工具，这些模型和工具涵盖了图像识别、自然语言处理等多个领域。
你的YOLO部署神器 TensorRT Plugin、CUDA Kernel、CUDA Graphs三管齐下，享受闪电般的推理速
最新发布
10-09
TensorRT-YOLO 是一个支持 YOLOv3、YOLOv5、YOLOv6、YOLOv7、YOLOv8、YOLOv9、YOLOv10、YOLO11、PP-YOLOE 和 PP-YOLOE+ 的推理加速项目，使用 NVIDIA TensorRT 进行优化。项目不仅集成了 TensorRT 插件以增强后处理效果，还使用了 CUDA 核函数以及 CUDA 图来加速推理。TensorRT-YOLO 提供了 C++ 和 Python 推理的支持，旨在提供快速而优化的目标检测解决方案。
Cuda Graph (cuda 优化)
qq_42105426的博客
02-08 1478
1、Nvidia 官网文档 cuda 官网文档 可查时间 2019年 https://developer.nvidia.com/blog/cuda-graphs/#:~:text=CUDA%20Graphs%20have%20been%20designed%20to%20allow%20work,a%20single%20CPU%20operation%2C%20and%20hence%20reduce%20overheads. 2、Cuda graph 作用 Cuda graph 通过掩盖 ker
CUDA Graph图详解
专注于人工智能领域的小何尚
05-31 7557
CUDA图 CUDA Graphs 为 CUDA 中的工作提交提供了一种新模型。图是一系列操作，例如内核启动，由依赖关系连接，独立于其执行定义。这允许一个图被定义一次，然后重复启动。将图的定义与其执行分开可以实现许多优化：首先，与流相比，CPU 启动成本降低，因为大部分设置都是提前完成的；其次，将整个工作流程呈现给 CUDA 可以实现优化，这可能无法通过流的分段工作提交机制实现。 要查看图形可能的优化，请考虑流中发生的情况：当您将内核放入流中时，主机驱动程序会执行一系列操作，以准备在 GPU 上执行内核。
CUDA效率优化之CUDA Graph的使用
二爷的博客
04-12 1878
GPU 架构的性能随着每一代的更新而不断提高。现代 GPU 每个操作（如kernel运行或内存复制）所花费的时间现在以微秒为单位。但是，将每个操作提交给 GPU 也会产生一些开销——也是微秒级的。实际的应用程序中经常要执行大量的 GPU 操作：典型模式涉及许多迭代（或时间步），每个步骤中有多个操作。如果这些操作中的每一个都单独提交到 GPU 启动并独立计算，那么提交启动开销汇总在一起可能导致明显的整体性能下降。CUDA Graphs 将整个计算流程定义为一个图而不是单个操作的列表。
在动态环境中使用CUDA图
专注于人工智能领域的小何尚
04-20 664
在动态环境中使用CUDA图 通过将计算密集型部分放到 GPU 上，可以大大加快许多工作负载的速度。 在 CUDA 术语中，这称为启动内核。 当这些内核很多且持续时间很短时，启动开销有时会成为问题。 CUDA Graphs 提供了一种减少这种开销的方法。 图之所以有效，是因为它们将任意数量的异步 CUDA API 调用（包括内核启动）组合成一个只需要一次启动的操作。 它们在创建时确实会产生一些开销，因此它们最大的好处来自于多次重用它们。 在工具包版本 10 中引入时，CUDA 图可能已经更新以反映其实例化的
CUDA中的图内存节点
专注于人工智能领域的小何尚
02-10 1132
图内存节点允许图创建和拥有内存分配功能。图内存节点具有 GPU 有序生命周期语义，它指示何时允许在设备上访问内存。这些 GPU 有序生命周期语义支持驱动程序管理的内存重用，并与流序分配 API和相匹配，这可能在创建图形时被捕获。图分配在图的生命周期内具有固定的地址，包括重复的实例化和启动。这允许图中的其他操作直接引用内存，而无需更新图，即使 CUDA 更改了后备物理内存也是如此。在一个图中，其图有序生命周期不重叠的分配可以使用相同的底层物理内存。
CUDA 版本对于深度学习任务性能的影响
CUDA（Compute Unified Device Architecture）是由 NVIDIA 推出的并行计算平台和编程模型。它允许开发者使用C/C++、Fortran等编程语言来访问GPU的高性能并行计算能力，加速应用程序的运行速度。CUDA包括一个GPU架构...
『大模型笔记』从基础原理出发提升深度学习性能
AI新视界
05-28 306
如果你想加速你的深度学习系统，最重要的是要弄清楚模型的瓶颈在哪里。这个瓶颈决定了你加速系统的最佳方法。我经常看到研究人员和其他想要加速PyTorch代码的人在不了解具体情况的前提下盲目尝试。性能环境可能的解决方案Overhead-Bound(开销瓶颈)Tracing, Operator Fusion, 不使用Python, 真正的JIT编译器Operator Fusion(算子融合)使用Tensor Cores, 购买更多Nvidia硬件当然，可以说，用户需要考虑这些问题本身就反映了框架的不足。
开源的深度学习库PyTorch介绍说明
02-17
1. **动态图（Dynamic Graphs）**： - **实时构建与修改**：与TensorFlow等其他框架采用的静态图不同，PyTorch使用动态图机制，允许用户在运行时动态地构建和修改计算图。这种特性对于快速实验和原型设计非常有利。...
ZeroShot Learning原理与代码实例讲解
AI天才研究院
07-28 357
在人工智能和机器学习领域中，模型的泛化能力一直是研究的重点。传统的监督学习方法通常需要大量的标注数据来训练模型，这不仅耗时耗力，而且在面对新类别或未见过的数据时往往表现不佳。为了解决这个问题，研究人员提出了零样本学习（ZeroShot Learning，简称ZSL）的概念。零样本学习是一种旨在识别或分类从未在训练集中出现过的类别的机器学习方法。它的核心思想是利用已知类别的语义信息来推断未知类别的特征，从而实现对新类别的识别。这种方法在实际应用中具有重要意义，特别是在快速变化的环境中或者数据收集困难的领域。
CUDA Graph（创建静态图加快后续多次启动速度）
smartcat2010的博客
03-28 6145
博客原文 应用场景：多次迭代（每次迭代执行的任务图都一样） 假设：shortKernel是一个运行时间极短的kernel 1. 初始版本： define NSTEP 1000 #define NKERNEL 20 // start CPU wallclock timer for(int istep=0; istep<NSTEP; istep++){ for(int ikrnl...
CUDA图形（Cuda Graphs）的实例编程
DevProZ的博客
09-15 339
通过使用CUDA图形，我们可以将多个CUDA操作组合成一个图形，并在GPU上执行，从而提高GPU应用程序的性能。CUDA图形的基本概念是将一系列CUDA操作打包成一个图形（Graph），然后在GPU上执行这个图形。下面是一个简单的示例，展示了如何使用CUDA图形来加速向量加法操作。通过使用CUDA图形，可以显著提高GPU应用程序的性能。本文将介绍如何在CUDA中使用图形来优化GPU计算，并提供相应的源代码。在上面的代码中，我们将设备上的输出向量复制到主机上的内存。是对应的CUDA核函数，
英伟达CUDA 10终于开放下载了
机器学习算法与Python学习
10-01 8360
栗子 发自 凹非寺量子位 出品 | 公众号 QbitAI英伟达的CUDA 10工具包，终于可以下载了。与9字头的前辈相比，10.0是第一次为图灵GPU定制的版本。英伟达新...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

Hi20240217

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值