CUDA流和事件

news/2024/9/24 20:42:49/

CUDA通过流来实现网格级并发。

流和事件

CUDA流是一系列异步的CUDA操作,这些操作按照主机代码确定的顺序在设备上执行。流可以封装这些操作,保持操作的顺序,允许操作在流中排队,并使他们在先前的所有操作之后执行。

这些操作包括在主机与设备之间进行数据传输,内核启动以及大多数由主机发起但由设备处理的其他命令。

流中操作的执行相对于主机总是异步的。CUDA运行时决定何时可以在设备上执行操作。我们的任务是使用CUDA的API来确保一个异步操作在运行结果被使用之前可以完成。

在同一个CUDA流中的操作有严格的执行顺序,而在不同的CUDA流中的操作在执行顺序上不受限制。使用多个流同时启动多个内核,可以实现网格级并发。

CUDA编程的一个典型模式是:

  1. 数据从主机移到设备上
  2. 设备上执行一个内核
  3. 将结果从设备移回主机

在多数情况下,执行内核比传输数据耗时的多。在这些情况下,可以完全隐藏CPU和GPU之间的通信延迟。将内核执行和数据传输调度到不同的流中,这些操作可以重叠,程序的总运行时间将被缩短。流在CUDA的API调用粒度上可实现流水线或双缓冲技术。

CUDA的api函数分为同步和异步。同步函数会阻塞主机端线程,直到其完成。异步函数被调用后,控制权直接移回主机。异步函数和流是在CUDA中构建网格级并发的两个基本支柱。

CUDA

所有的CUDA操作(内核和数据传输)都在一个流中显式或隐式的运行。流分为两种:

  • 隐式声明的流(空流)
  • 显式声明的流(非空流)

如果没有显式的声明一个流,那么内核启动和数据传输将默认使用空流。

非空流可以被显式的创建和管理。如果想要重叠不同的CUDA操作,必须使用非空流。基于流的异步的内核启动和数据传输支持以下类型的并发:

  • 重叠主机计算和设备计算
  • 重叠主机计算和主机与设备间的数据传输
  • 重叠主机与设备间的数据传输和设备计算
  • 并发设备计算

思考下面使用默认流的方法:

cudaMemcpy(...,cudaMemcpyHostToDevice);
kernel<<<grid,block>>>();
cudaMemcpy(...,cudaMemcpyDeviceToHost);

从设备角度来看,上述代码的所有3个操作都被发布到默认流中,并且按发布顺序执行。设备不知道其他被执行的主机操作。

从主机角度看,数据传输是同步的,强制空闲主机等待数据传输完成。内核启动是异步的,无论内核是否完成,主机的应用程序都立即恢复执行。这种内核启动的默认异步行为使它可以直接重叠设备和主机计算

数据传输也可以异步发布,但是必须显式的设置一个CUDA流来装载。提供以下函数

cudaError_t cudaMemcpyAsync(void* dst, void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream=0);

附加的流的标识符作为第五个参数,默认情况下使用默认流。这个函数与主机是异步的,所以调用发布后,控制权将立即返回到主机。

使用如下代码创建一个非空流

cudaError_t cudaStreamCreate(cudaStream_t* pStream);

该函数创建了一个可以显示管理的非空流。之后,返回到pStream中的流就可以被当做流参数供其他的异步CUDA的API函数使用。

在执行异步数据传输时,必须使用固定主机内存。在非默认流中启动内核,必须在内核执行配置中提供一个流标识符作为第四个参数

kernel<<<grid, block, shareMemSize, stream>>>()

非默认流的声明和创建如下:

cudaStream_t stream;
cudaStreamCreate(&stream);//释放资源
cudaStreamDestroy(cudaStream_t stream);

在一个流中,当cudaStreamDestroy被调用时,如果该流中仍有未完成的工作,函数将立即返回,当流中所有工作都已完成时,与流相关的资源将被自动释放。

用以下两个函数检查流中工作是否完成

cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);

synchronize将阻塞主机,直到给定流中的工作全部完成。query检查流中所有操作是否完成,但不会阻塞,如果完成返回cudaSuccess,未完成返回cudaErrorNotReady。

画个图展示流的作用

for(int i=0; i<nStreams; i++){int offset = i * bytesPerStream;cudaMemcpyAsync(&d_a[offset], &a[offset], bytesPerStream, streams[i]);kernel<<<grid, block, 0, streams[i]>>>(&d_a[offset]);cudaMemcpyAsync(&a[offset], &d_a[offset], bytesPerStream, streams[i]);
}for(int i=0; i<nStreams; i++){cudaStreamSynchronize(streams[i]);
}

并发内核的最大数量是依赖设备的。费米架构支持16路并发,开普勒设备支持32路并发。

流调度

从概念上讲,所有流可以并发执行。但是,当流映射到物理硬件时,并不总是这样的。

虚假的依赖关系

虚假的依赖关系(False Dependency)指的是一种由编译器或硬件引入的假象性依赖,导致了代码中不必要的序列化或延迟。这种依赖关系并不反映真实的数据依赖关系,但会影响到代码的执行顺序和性能。

虚假的依赖关系通常出现在对共享内存的操作中,特别是在使用指针进行多次访存的情况下。编译器或硬件可能会认为对同一内存地址的多次访问之间存在依赖关系,从而引入不必要的序列化或延迟。

Hyper-Q技术

Hyper-Q技术使用多个硬件工作队列,减少了虚假的依赖关系。允许多个CPU线程或进程在单一GPU上同时启动工作。

开普勒GPU使用32个工作队列,每个流分配一个工作队列。如果超过32个流,多个流将共享一个硬件工作队列。

流的优先级

计算能力3.5或更高的设备,可以给流分配优先级。优先级高的流的网格队列可以优先占有低优先级流的已经执行的工作。

CUDA事件

CUDA中事件本质上是CUDA流中的标记,它与流内操作中特定点相关联。可以使用事件来执行以下两个基本任务:

  • 同步流的执行
  • 监控设备的进展

CUDA的API提供了在流中任意点插入事件以及查询事件完成的函数。只有当一个给定CUDA流中先前的所有操作都执行结束后,记录在该流内的事件才会起作用(即完成)。

创建和销毁

声明

cudaEvent_t event;

创建

cudaError_t cudaEventCreate(cudaEvent_t* event);

销毁

cudaError_t cudaEventDestroy(cudaEvent event);

记录事件和计算运行时间

事件在流中标记了一个点。可以用来检查正在执行的流操作是否已到达了给定点。

使用如下函数进入CUDA

cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream=0);

使用如下函数计算两个事件标记的CUDA操作运行时间

cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t end);

返回运行时间,单位为ms。

示例如何进行计时

//创建两个事件
cudaEvent_t start, end;
cudaEventCreate(&start);
cudaEventCreate(&end);//将开始事件记录在默认流中
cudaEventRecord(start);kernel<<<>>>();//将结束事件记录在默认流中
cudaEventRecord(end);//等待事件结束
cudaEventSynchronize(end);//计算时间
float time;
cudaEventElapsedTime(&time, start, end);//释放事件
cudaEventDestroy(start);
cudaEventDestroy(end);

cudaEventSynchronize和stream的相关函数相同。

CUDA流中的事件可以在主机端和设备端都记录。在CUDA中,事件(Event)用于测量时间间隔或同步CUDA流中的操作。主要有以下两种类型的事件:

  1. 主机事件(Host Event):主机事件是由主机代码创建和记录的事件,用于测量主机和设备之间的时间间隔或同步主机代码和CUDA流中的操作。可以使用cudaEventRecord()函数记录主机事件。

  2. 设备事件(Device Event):设备事件是由设备代码(即在CUDA核函数中)创建和记录的事件,用于测量CUDA流中的操作的时间间隔或同步不同的CUDA核函数。可以使用cudaEventRecord()函数在设备代码中记录设备事件。

在使用CUDA流时,通常会在主机端记录主机事件来测量主机与设备之间的时间间隔,同时也可以在设备端使用设备事件来测量CUDA核函数的执行时间或同步不同核函数之间的操作。


http://www.ppmy.cn/news/1454705.html

相关文章

7 - 函数式编程

文章目录 Interface函数式接口Lambda表达式语法格式替代匿名类集合迭代方法引用作用域范围this关键字 Stream流流类型创建流中间操作流过滤 filter映射 map去重 distinct排序 sorted限制 limit跳过 skipflatMap 终结操作流forEachcountmin/maxcollect匹配组合 reduce 延迟执行并…

Py深度学习基础|python中类的特殊方法-__getitem__()

1.基本介绍 在Python中&#xff0c;__getitem__是一个特殊方法&#xff08;也常被称为“魔术方法”&#xff0c;即双下划线方法&#xff09;&#xff0c;它使一个类的实例对象能够支持通过键来获取其内部数据&#xff0c;类似于操作列表、元组或字典的方式。当你尝试使用方括号…

OpenGL 的内置矩阵种种

文章目录 一、说明二、glMatrixMode的三个选项三、OpenGL 矩阵操作四、入栈和弹出矩阵五、设置内置 OpenGL 矩阵六、矩阵设置异端七、正常矩阵 一、说明 内置 OpenGL有几个重要矩阵&#xff0c;围绕这几个矩阵&#xff0c;OpenGL 有一小组 矩阵操作。在本例中为 glMatrixMode&a…

RFC 791 (1)-导论

目录 浅论 IP是啥 IP可以管啥 操作 范例查看 提示&#xff1a;本系列将会开始RFC文档阅读&#xff0c;这里会给出我的一些笔记 浅论 我们这篇RFC文档描述的是IP和ICMP协议&#xff0c;我们都知道&#xff0c;在传统的OSI七层或者是现在被简化的五层&#xff1a;应用层&…

Docker 入门与实践:从零开始构建容器化应用环境

Docker 一、docker常用命令docker ps 格式化输出Linux设置命令别名 二、数据卷相关命令挂载到默认目录&#xff08;/var/lib/docker&#xff09;挂载到本地目录 三、自定义镜像Dockerfile构建镜像的命令 四、网络自定义网络 五、DockerCompose相关命令 一、docker常用命令 dock…

知识图谱实用网站推荐

​ 知识图谱是一种用于描述和组织的图形化表示方法&#xff0c;它是一种基于语义的知识表示方式&#xff0c;可以将各种实体、概念、属性以及关系等知识元素以图谱的形式进行展示和组织。 本文给大家分享了如何实现知识图谱中图关系的构建&#xff0c;以下是相关网站&#xff…

从新手到聊天专家:ChatGPT对话技巧全解析

关于为什么AI未能为你所用&#xff0c;许多人在拿到GPT后的首个问题便是&#xff1a;使用不佳。 这意味着&#xff0c;你可能会觉得ChatGPT的回答过于空洞&#xff0c;缺乏足够的参考价值。 其次&#xff0c;一个常见的问题是&#xff1a;不知如何利用。 即便拥有了GPT&#…

OpenHarmony实战开发-动画曲线、如何实现动画衔接

UI界面除了运行动画之外&#xff0c;还承载着与用户进行实时交互的功能。当用户行为根据意图变化发生改变时&#xff0c;UI界面应做到即时响应。例如用户在应用启动过程中&#xff0c;上滑退出&#xff0c;那么启动动画应该立即过渡到退出动画&#xff0c;而不应该等启动动画完…