0
  • 聊天消息
  • 系统消息
  • 评论与回复
登录后你可以
  • 下载海量资料
  • 学习在线课程
  • 观看技术视频
  • 写文章/发帖/加入社区
会员中心
创作中心

完善资料让更多小伙伴认识你,还能领取20积分哦,立即完善>

3天内不再提示

如何在CUDA C/C++中实现数据传输和其他操作的重叠

星星科技指导员 来源:NVIDIA 作者:Mark Harris 2022-04-11 10:37 次阅读

在上一期的 C / C ++ 文章 中,我们讨论了如何在主机和设备之间高效地传输数据。在这篇文章中,我们讨论了如何将数据传输与主机上的计算、设备上的计算相重叠,在某些情况下,主机和设备之间的其他数据传输。实现数据传输和其他操作之间的重叠需要使用 CUDA 流,所以首先让我们了解一下流。

CUDA 流

CUDA 中的 stream 是按照主机代码发出的顺序在设备上执行的操作序列。虽然流中的操作被保证按规定的顺序执行,但是不同流中的操作可以被交错,并且在可能的情况下,它们甚至可以并发运行。

默认流

CUDA 中的所有设备操作(内核和数据传输)都在一个流中运行。如果没有指定流,则使用默认流(也称为“空流”)。默认流与其他流不同,因为它是关于设备上操作的同步流:在所有先前发出的操作 在设备上的任何流中 完成之前,默认流中的任何操作都不会开始,并且默认流中的操作必须在任何其他操作(在设备上的任何流中)之前完成就要开始了。

请注意, 2015 年发布的 CUDA 7 引入了一个新的选项,即每个主机线程使用单独的默认流,并将每个线程的默认流视为常规流(即它们不与其他流中的操作同步)。在文章 GPU 专业提示: CUDA 7 流简化并发 中阅读更多关于这种新行为的信息

让我们看一些使用默认流的简单代码示例,并从主机和设备的角度讨论操作是如何进行的。

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a)
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

在上面的代码中,从设备的角度来看,所有三个操作都被发布到同一个(默认)流中,并将按照它们发出的顺序执行。

从主机的角度看,隐式数据传输是阻塞或同步传输,而内核启动是异步的。由于第一行上的主机到设备的数据传输是同步的, CPU 线程在主机到设备的传输完成之前不会到达第二行的内核调用。一旦内核被发出, CPU 线程将移动到第三行,但由于设备端的执行顺序,该行上的传输无法开始。

内核从主机的角度启动的异步行为使得重叠的设备和主机计算非常简单。我们可以修改代码以添加一些独立的 CPU 计算,如下所示。

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a)
myCpuFunction(b)
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

在上面的代码中,一旦 increment() 内核在设备上启动, CPU 线程就执行 myCpuFunction() ,它在 CPU 上的执行与在 GPU 上的内核执行重叠。无论是主机功能还是设备内核先完成,都不会影响后续的设备到主机的传输,只有在内核完成后才会开始,从设备的角度来看,上一个例子没有什么变化,设备完全不知道 myCpuFunction() 。

非默认流

在下面的代码中, CUDA C / C ++的非默认流被声明、创建和销毁。

cudaStream_t stream1;
cudaError_t result;
result = cudaStreamCreate(&stream1)
result = cudaStreamDestroy(stream1)

为了向非默认流发出数据传输,我们使用了cudaMemcpyAsync()函数,它类似于前一篇文章中讨论的cudaMemcpy()函数,但将流标识符作为第五个参数

result = cudaMemcpyAsync(d_a, a, N, cudaMemcpyHostToDevice, stream1)

cudaMemcpyAsync() 在主机上是非阻塞的,因此在发出传输之后,控制权立即返回到主机线程。此例程有 cudaMemcpy2DAsync() 和 cudaMemcpy3DAsync() 变体,它们可以在指定的流中异步传输 2D 和 3D 数组部分。

为了向非默认流发出内核,我们将流标识符指定为第四个执行配置参数(第三个执行配置参数分配共享设备内存,我们将在后面讨论;现在使用 0 )。

increment<<<1,N,0,stream1>>>(d_a)

与流同步

由于非默认流中的所有操作相对于宿主代码都是非阻塞的,因此您将遇到需要将宿主代码与流中的操作同步的情况。“重锤”的方法是使用 cudaDeviceSynchronize() ,它会阻止主机代码,直到之前在设备上发出的所有操作都完成为止。在大多数情况下,这是一种过度杀戮,并且会由于整个设备和主机线程的暂停而影响性能。

CUDA 流 API 有多种不太严格的同步主机与流的方法。函数 cudaStreamSynchronize(stream) 可用于阻止主机线程,直到指定流中以前发出的所有操作都已完成。函数 cudaStreamQuery(stream) 测试向指定流发出的所有操作是否已完成,而不阻止主机执行。函数 cudaEventSynchronize(event) 和 cudaEventQuery(event) 的行为与它们的流对应项相似,只是它们的结果基于是否记录了指定的事件,而不是基于指定的流是否空闲。您还可以使用 cudaStreamWaitEvent ( event )在单个流中同步特定事件的操作(即使事件记录在不同的流中,或者记录在不同的设备上)。

重叠的内核执行和数据传输

前面我们演示了如何将默认流中的内核执行与主机上的代码执行重叠。但我们在这篇文章中的主要目标是向您展示如何将内核执行与数据传输重叠。要做到这一点有几个要求。

设备必须能够“并发复制和执行”。这可以从 cudaDeviceProp 结构的 deviceOverlap 字段或从 CUDA SDK / Toolkit 附带的 deviceQuery 示例的输出中进行查询。几乎所有具有计算能力 1 。 1 及更高版本的设备都具有此功能。

要重叠的内核执行和数据传输必须同时发生在 different 、 non-default 流中。

数据传输所涉及的主机内存必须是 pinned 内存。

因此,让我们从上面修改我们的简单主机代码,以使用多个流,看看是否可以实现任何重叠。这个例子的完整代码是 在 Github 上提供 。在修改后的代码中,我们将大小为 N 的数组分解为 streamSize 元素的块。由于内核对所有元素都是独立操作的,因此每个块都可以独立处理。使用的(非默认)流数为 nStreams=N/streamSize 。有多种方法可以实现数据的域分解和处理;一种方法是循环使用数组中每个块的所有操作,如本示例代码所示。

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  cudaMemcpyAsync(&d_a[offset], &a[offset], streamBytes, cudaMemcpyHostToDevice, stream[i]);
  kernel<<>>(d_a, offset);
  cudaMemcpyAsync(&a[offset], &d_a[offset], streamBytes, cudaMemcpyDeviceToHost, stream[i]);
}

另一种方法是将类似的操作批处理在一起,首先发出所有主机到设备的传输,然后是所有的内核启动,然后是所有设备到主机的传输,如下面的代码所示。

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  cudaMemcpyAsync(&d_a[offset], &a[offset],
                  streamBytes, cudaMemcpyHostToDevice, cudaMemcpyHostToDevice, stream[i]);
}

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  kernel<<>>(d_a, offset);
}

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  cudaMemcpyAsync(&a[offset], &d_a[offset],
                  streamBytes, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToHost, stream[i]);
}

上面显示的两个异步方法都会产生正确的结果,并且在这两种情况下,依赖操作都会按照它们需要执行的顺序发布到同一个流。但根据所使用的 GPU 的特定代数,这两种方法的性能截然不同。在 Tesla C1060 (计算能力 1 。 3 )上运行测试代码(来自 Github )给出以下结果。

Device : Tesla C1060

Time for sequential transfer and execute (ms ): 12.92381
  max error : 2.3841858E -07
Time for asynchronous V1 transfer and execute (ms ): 13.63690
  max error : 2.3841858E -07
Time for asynchronous V2 transfer and execute (ms ): 8.84588
  max error : 2.3841858E -07

在 Tesla C2050 (计算能力 2 . 0 )上,我们得到以下结果。

Device : Tesla C2050

Time for sequential transfer and execute (ms ): 9.984512
  max error : 1.1920929e -07
Time for asynchronous V1 transfer and execute (ms ): 5.735584
  max error : 1.1920929e -07
Time for asynchronous V2 transfer and execute (ms ): 7.597984
  max error : 1.1920929e -07

这里第一次报告的是使用阻塞传输的顺序传输和内核执行,我们将其作为异步加速比较的基线。为什么这两种异步策略在不同的体系结构上表现不同?要破解这些结果,我们需要更多地了解 CUDA 设备如何调度和执行任务。 CUDA 设备包含用于各种任务的引擎,这些引擎在发出操作时对操作进行排队。不同引擎中的任务之间的依赖关系得到维护,但是在任何引擎中,所有外部依赖关系都会丢失;每个引擎队列中的任务将按照它们的发出顺序执行。 C1060 有一个拷贝引擎和一个内核引擎。在 C1060 上执行示例代码的时间线如下图所示。

在这个示意图中,我们假设主机到设备传输、内核执行和设备到主机传输所需的时间大致相同(选择内核代码是为了实现这一点)。正如顺序内核所期望的那样,任何操作中都没有重叠。对于我们代码的第一个异步版本,复制引擎中的执行顺序是: H2D stream ( 1 )、 D2H stream ( 1 )、 H2D stream ( 2 )、 D2H stream ( 2 )等等。这就是为什么我们在 C1060 上使用第一个异步版本时看不到任何加速:任务是按照排除内核执行和数据传输重叠的顺序被发送到复制引擎的。然而,对于版本 2 ,在所有主机到设备的传输在任何设备到主机的传输之前发出,重叠是可能的,如较低的执行时间所示。根据我们的示意图,我们期望异步版本 2 的执行时间是顺序版本的 8 / 12 ,或者 8 。 7ms ,这在前面给出的计时结果中得到了确认。

在 C2050 上,两个功能相互作用导致与 C1060 不同的行为。 C2050 有两个复制引擎,一个用于主机到设备的传输,另一个用于设备到主机的传输,以及一个内核引擎。下图说明了我们的示例在 C2050 上的执行。

有两个复制引擎解释了为什么异步版本 1 在 C2050 上实现了很好的加速:流[i] 不阻止流中数据的主机到设备传输 [i + 1]中数据的主机到设备的传输,因为 C2050 上的每个复制方向都有一个单独的引擎。示意图预测了执行情况相对于顺序版本,时间被缩短一半,这大致就是我们的计时结果显示的。

但是在 C2050 上的异步版本 2 中观察到的性能下降呢?这与 C2050 并发运行多个内核的能力有关。当多个内核在不同(非默认)流中背靠背地发出时,调度程序尝试启用这些内核的并发执行,结果会延迟通常在每个内核完成后出现的信号(这负责启动设备到主机的传输),直到所有内核完成。因此,虽然在第二个版本的异步代码中,主机到设备的传输和内核的执行之间有重叠,但是内核执行和设备到主机的传输之间没有重叠。示意图预测异步版本 2 的总时间是顺序版本的 9 / 12 ,即 7 。 5 毫秒,这一点由我们的计时结果证实。

CUDA Fortran 异步数据传输 中提供了关于本文中使用的示例的更详细的描述,好消息是对于具有计算能力 3 。 5 ( K20 系列)的设备, Hyper-Q 特性消除了定制发布顺序的需要,因此上述任何一种方法都可以工作。我们将在以后的文章中讨论使用开普勒特性,但是现在,这里是在 Tesla K20c GPU 上运行示例代码的结果。如您所见,这两个异步方法在同步代码上实现了相同的加速。

Device : Tesla K20c
Time for sequential transfer and execute (ms): 7.101760
  max error : 1.1920929e -07
Time for asynchronous V1 transfer and execute (ms): 3.974144
  max error : 1.1920929e -07
Time for asynchronous V2 transfer and execute (ms): 3.967616
  max error : 1.1920929e -07

概括

这篇文章和 上一个 讨论了如何优化主机和设备之间的数据传输。上一篇文章集中讨论了如何最小化执行这种传输的时间,这篇文章介绍了流,以及如何使用流通过并发执行副本和内核来屏蔽数据传输时间。

在一篇关于流的文章中,我应该提到,虽然使用默认流可以方便地开发代码,但同步代码更简单,最终您的代码应该使用非默认流或 CUDA 7 对每线程默认流的支持(读 GPU 专业提示: CUDA 7 流简化并发 )。这在编写库时尤其重要。如果库中的代码使用默认流,那么最终用户就没有机会将数据传输与库内核执行重叠。

现在您已经知道如何在主机和设备之间高效地移动数据,所以我们将研究如何在 下一篇文章 中的内核中高效地访问数据。

关于作者

Mark Harris 是 NVIDIA 杰出的工程师,致力于 RAPIDS 。 Mark 拥有超过 20 年的 GPUs 软件开发经验,从图形和游戏到基于物理的模拟,到并行算法和高性能计算。当他还是北卡罗来纳大学的博士生时,他意识到了一种新生的趋势,并为此创造了一个名字: GPGPU (图形处理单元上的通用计算)。

审核编辑:郭婷

声明:本文内容及配图由入驻作者撰写或者入驻合作网站授权转载。文章观点仅代表作者本人,不代表电子发烧友网立场。文章及其配图仅供工程师学习之用,如有内容侵权或者其他违规问题,请联系本站处理。 举报投诉
  • NVIDIA
    +关注

    关注

    14

    文章

    4862

    浏览量

    102731
  • gpu
    gpu
    +关注

    关注

    28

    文章

    4682

    浏览量

    128623
  • C++
    C++
    +关注

    关注

    21

    文章

    2102

    浏览量

    73460
收藏 人收藏

    评论

    相关推荐

    socket 数据传输效率提升技巧

    在现代网络应用数据传输效率是衡量系统性能的关键指标之一。对于使用socket进行数据传输的应用,优化传输效率不仅可以提升用户体验,还能降低成本。 1. 选择合适的
    的头像 发表于 11-12 14:34 98次阅读

    CAN总线数据传输速率设置

    CAN(Controller Area Network)总线是一种串行通信协议,主要用于汽车和工业控制系统,以实现电子控制单元(ECU)之间的通信。CAN总线的数据传输速率,也称为波特率,是衡量
    的头像 发表于 11-12 10:03 88次阅读

    LORA模块的数据传输速率

    有所不同。以下是关于LoRa模块数据传输速率的一些关键点: 数据传输速率的可变性 : LoRa技术允许在不同的数据速率下操作,以适应不同的应用需求。速率可以从几百比特每秒(bps)到几
    的头像 发表于 10-31 17:03 530次阅读

    网络数据传输速率的单位是什么

    网络数据传输速率的单位是 bps(bit per second) ,即比特每秒,也可以表示为b/s或bit/s。它表示的是每秒钟传输的二进制数的位数。比特(bit)是计算机数据量的单
    的头像 发表于 10-12 10:20 814次阅读

    高速串行总线,数据传输离不开它!#高速串行总线 #电路知识 #数据传输

    电路数据传输
    安泰仪器维修
    发布于 :2024年08月20日 15:42:00

    C++实现类似instanceof的方法

    函数,可实际上C++没有。但是别着急,其实C++中有两种简单的方法可以实现类似Java的instanceof的功能。 在
    的头像 发表于 07-18 10:16 520次阅读
    <b class='flag-5'>C++</b><b class='flag-5'>中</b><b class='flag-5'>实现</b>类似instanceof的方法

    边OTG边充电芯片如何实现充电与数据传输并行?

    边OTG边充电芯片实现充电与数据传输并行的功能,主要依赖于其内部的设计和与USB Type-C接口标准的结合。
    的头像 发表于 07-14 10:35 488次阅读

    以太网接口的数据传输原理详解

    以太网接口作为计算机网络的关键组成部分,承担着数据传输的重要职责。在了解以太网接口的数据传输原理之前,我们首先需要明确以太网的基本概念和工作机制。以太网是一种广泛应用的局域网技术,它基于CSMA
    的头像 发表于 05-29 16:47 1441次阅读

    探索SPI单线传输模式时钟线与数据传输的简化

    通信的简化需求也日益增加。在这种背景下,SPI的单线传输模式成为了一个备受关注的解决方案。 SPI协议概述 SPI协议是一种常用的同步串行通信协议,通常用于微控制器与其他设备之间的数据传输。它基于主从架构,允许一个主机与多个从机
    的头像 发表于 05-28 18:26 1071次阅读

    GMSL技术 实现高带宽、低延迟和高可靠性数据传输# ADI# GMSL# 汽车# 数据传输

    adi数据传输电机
    Excelpoint世健
    发布于 :2024年05月17日 16:34:25

    DTU的多种协议,解锁数据传输的无限可能

    DTU,即数据传输单元,是一种在物联网(IoT)网络中常用的设备,主要用于在传感器和智能设备之间进行数据传输。DTU使用多种协议来实现这一目标,这些协议不仅提高了数据传输的效率,还增强
    的头像 发表于 03-01 11:00 720次阅读
    DTU的多种协议,解锁<b class='flag-5'>数据传输</b>的无限可能

    两路数据传输,CY7C68013都作为从机接收数据是否可行?

    我想使用CY7C68013A的GPIF和FIFO功能: 1. 两路数据传输,CY7C68013都作为从机接收数据 2. 每一路都数据格式为:
    发表于 02-28 07:34

    手机没有OTG功能,如何实现数据传输

    手机没有OTG功能,如何实现数据传输? 手机没有OTG功能,需要传输数据的时候可以考虑以下几种方法: 1. 云端存储 云端存储是目前非常流行的一种
    的头像 发表于 12-11 15:31 3237次阅读

    WT2605C-L011语音芯片IC:蓝牙音频与数据传输功能的集成优势及应用

    随着无线技术的飞速发展,蓝牙技术在音频和数据传输领域的应用越来越广泛。作为一款领先的语音芯片,WT2605C-L011不仅支持蓝牙音频播放,还具备蓝牙BLE数据传输功能。同时,其支持
    的头像 发表于 11-27 10:07 574次阅读
    WT2605<b class='flag-5'>C</b>-L011语音芯片IC:蓝牙音频与<b class='flag-5'>数据传输</b>功能的集成优势及应用

    芯片设计半双工和全双工数据传输的区别

    在现代通信技术,半双工和全双工数据传输是两种常见的数据传输方式。本文将为大家详细解析这两种传输方式在芯片设计的应用和区别,帮助大家更好地
    发表于 11-19 11:31 862次阅读