当前位置： 首页 > news >正文

CUDA Graphs学习与实验

news 2025/8/18 6:51:58

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

一.参考链接

  • graph management functions of the low-level CUDA driver api
  • CUDA Runtime Graph API

二.测试方案

请添加图片描述

三.测试代码

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
HostFunc13.0013.0013.0013.0013.0013.0013.0013.00
graph mode
Kernel1
Kernel2
Kernel3
Kernel4
Kernel5
HostFunc13.0013.0013.0013.0013.0013.0013.0013.00

请添加图片描述

http://www.lryc.cn/news/457702.html

相关文章：

  • 【自注意力与Transformer架构在自然语言处理中的演变与应用】
  • LabVIEW交直流接触器动态检测系统
  • Unity3D中基于四叉树的范围检测算法详解
  • k8s网络通信
  • 07 欢乐的跳
  • 【韩顺平Java笔记】第8章：面向对象编程（中级部分）【262-271】
  • GNU链接器（LD）：输入分区的垃圾回收及保护处理（KEEP命令）介绍
  • 论文翻译 | Fairness-guided Few-shot Prompting for LargeLanguage Models
  • 【分布式微服务云原生】战胜Redis脑裂：深入解析与解决方案
  • 数据治理与可持续发展：开启企业价值新模式——The Open Group 2024生态系统架构·可持续发展年度大会邀您共襄盛举
  • 数据库的分类及主流数据库
  • Qt C++设计模式-＞备忘录模式
  • Vue使用@别名替换后端ip地址
  • 强大的PDF到Word转换工具
  • js进阶——深入解析JavaScript中的URLSearchParams
  • 如何利用wsl-Ubuntu里conda用来给Windows的PyCharm开发
  • 操作系统的了解及安装
  • 【C++篇】虚境探微：多态的流动诗篇，解锁动态的艺术密码
  • uniapp的相关知识（1）
  • uniapp生成随机数
  • 使用jenkins将airflow-dbt部署到服务器上
  • 初学java练习题【1】
  • 大模型应用探讨，免费AI写作、一键PPT、免费PDF百种应用、与AI对话
  • 计算机视觉之OpenCV vs YOLO
  • 【深度学习基础模型】胶囊网络（Capsule Networks, CapsNet）详细理解并附实现代码。
  • 科普向 -- 什么是RPC
  • SpringBoot教程（二十四） | SpringBoot实现分布式定时任务之Quartz（基础）
  • 【现代控制理论】第2-5章课后题刷题笔记
  • （四）Proteus仿真STM32单片机使用定时器控制LED
  • Python快速编程小案例——打印蚂蚁森林植树证书