CUDA编程第六章: 流和并发

最新推荐文章于 2023-08-23 21:30:30 发布
最新推荐文章于 2023-08-23 21:30:30 发布
阅读量2.9k 收藏 28
点赞数 7
分类专栏： CUDA
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/qq_42683011/article/details/113829125
版权

本章内容：

  • 理解流和事件的本质

  • 理解网格级并发

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

  • 重叠CPU和GPU执行

  • 理解同步机制

  • 避免不必要的同步

  • 调整流的优先级

  • 注册设备回调函数

  • 通过NVIDIA可视化性能分析器显示应用程序执行的时间轴

一般来说，在CUDA C编程中有两个级别的并发：

  • 内核级并发

  • 网格级并发

到目前为止，你的关注点可能仅限于内核级的并发，在此级别的并发中，单一的任务或内核被GPU的多个线程并行执行。前面几章已经介绍了提升内核性能的几种方法，它们分别是从编程模型、执行模型和内存模型的角度进行介绍的。想必你已经了解了一些通过命令行性能分析器来研究和分析内核行为的方法。

本章将研究网格级的并发。在网格级并发中，多个内核在同一设备上同时执行，这往往会让设备利用率更好。在本章中，你将学习到如何使用CUDA流实现网格级的并发。还将使用CUDA的可视化性能分析器nvvp将内核并发执行可视化

6.1 流和事件概述:

CUDA流是一系列异步的CUDA操作，这些操作按照主机代码确定的顺序在设备上执行

CUDA流更像是一个由操作组成的队列, 流中的操作按照指定的顺序执行, 与主机保持异步状态, 并可查询队列状态

流中的操作是按照严格的顺序执行的, 但是不同的流中的操作执行顺序则不受限制
所以使用多个流同时启动多个内核, 可以实现网格级的并发, 并且能隐藏这些操作的延时

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

CUDA流:

所有的CUDA操作（包括内核和数据传输）都在一个流中显式或隐式地运行

流分为两种类型：

  • 隐式声明的流（空流）

  • 显式声明的流（非空流）

如果没有显式地指定一个流，那么内核启动和数据传输将默认使用空流
所以之前使用的都是空流或默认流

另一方面, 非空流可以被显式地创建和管理
并且, 如果想要重叠不同的CUDA操作，必须使用非空流

基于流的异步的内核启动和数据传输支持以下类型的粗粒度并发：

  • 重叠主机计算和设备计算

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

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

  • 并发设备计算

参考之前的默认流代码:

image-20210216115150461

  • 两个Memcpy是同步的
  • 核函数调用是异步的

流的创建

流可以将Memcpy操作也变成异步的:

首先创建一个非空流:

__host__ cudaError_t cudaStreamCreate(cudaStream_t * pStream)

而后使用异步的Memcpy函数:

这里的函数名就是多了个Async, 翻译为异步

__host__ cudaError_t cudaMemcpyAsync(void * dst, const void * src, size_t count, cudaMemcpyKind kind, cudaStream_t stream __dv (0))

参数解释:

  • 前头的四个参数与普通的Memcpy相同
  • 第五个参数则是指定使用的流
    为空的情况则使用默认流 (标识符为0)

这里需要额外注意的是, 使用异步数据传输时, 必须使用固定(非分页)的主机内存

image-20210216115949675

流的使用

在核函数的启动上, 也与使用默认流不同:

需要提供一个额外的流标识符作为第四个参数

image-20210216120108657

流的销毁

使用完流之后, 也可以手动销毁:

image-20210216120427701

当流中的工作尚未完成时, 销毁失败, cudaStreamDestroy函数将立刻返回

只有当流内的操作全部完成时, 才会正常销毁

流的查询:

因为所有的CUDA流操作都是异步的，所以CUDA的API提供了两个函数来检查流中所有操作是否都已经完成

image-20210216120822981

  • cudaStreamSynchronize强制阻塞主机，直到在给定流中所有的操作都完成了

  • cuda-StreamQuery会检查流中所有操作是否都已经完成，但在它们完成前不会阻塞主机:

    当所有操作都完成时cudaStreamQuery函数会返回cudaSuccess
    当一个或多个操作仍在执行或等待执行时返回cudaErrorNotReady

一个简单的栗子:

以下展现了在多个流中调度CUDA操作的栗子

image-20210216120910959

image-20210216121000063

在图6-1中，数据传输操作虽然分布在不同的流中，但是并没有并发执行。这是由一个共享资源导致的：PCIe总线。虽然从编程模型的角度来看这些操作是独立的，但是因为它们共享一个相同的硬件资源，所以它们的执行必须是串行的。具有双工PCIe总线的设备可以重叠两个数据传输，但它们必须在不同的流中以及不同的方向上。在图6-1中可以观察到，在一个流中从主机到设备的数据传输与另一个流中从设备到主机的数据传输是重叠的

并发内核的最大数量是依赖设备而确定的

Fermi设备支持16路并发，Kepler设备支持32路并发。设备上可用的计算资源进一步限制了并发内核的数量，如共享内存和寄存器。在本章后面的例子中将会探索这些局限性

流调度:

从逻辑上讲, 所有的流都可以并发执行

但是受限于物理硬件, 流需要合理的调度才能正确的执行

虚假的依赖关系:

这里需要介绍CUDA运行时的一个特点:

工作队列中，一个被阻塞的操作会将队列中该操作后面的所有操作都阻塞，即使它们属于不同的流
这种不同流的任务之间的依赖关系就是虚假依赖关系

如图所示:

image-20210216160446667

在流的工作队列中, 后续工作需要等待前头的工作完成之后才能进行, 这种就是依赖关系(有点类似工序图)

所以如果按照图中输入的三个流, 仅有画圈的部分能够被并行, 其余的都是串行执行

Hyper-Q技术:

Kepler架构中引入的Hyper-Q技术

其使用多个硬件工作队列，从而减少了虚假的依赖关系
Hyper-Q技术通过在主机和设备之间维持多个硬件管理上的连接，允许多个CPU线程或进程在单一GPU上同时启动工作

被Fermi架构中虚假依赖关系限制的应用程序，在不改变任何现有代码的情况下可以看到显著的性能提升

image-20210216160917577

但需要注意的是, Kepler GPU使用32个硬件工作队列, 每个流分配一个工作队列
如果启动的流熟练超过32个, 则将出现多个流共享一个工作队列的情况, 增加虚假依赖关系

流的优先级:

对于计算能力>=3.5的设备, 支持给流分配优先级:

__host__ cudaError_t cudaStreamCreateWithPriority(cudaStream_t * pStream, unsigned int flags, int priority)

这个函数创建了一个具有指定整数优先级的流，并在pStream中返回一个句柄

如果优先级设定超过了允许范围, 将被自动调整为min或max

可以使用此函数查询优先级的允许范围:

__host__ __cudart_builtin__ cudaDeviceGetStreamPriorityRange(int * leastPriority, int * greatestPriority)

这个函数的返回值存放在leastPriority和greatestPriority中，分别对应于当前设备的最低和最高优先级

当设备不支持优先级时, 将返回两个0

流优先级不会影响数据传输操作，只对计算内核有影响
高优先级流的网格队列可以优先占有低优先级流已经执行的工作

CUDA计算能力&显卡对照表:

最新信息参考这个:

https://developer.nvidia.com/cuda-gpus

Tesla Workstation Products
GPUCompute Capability
Tesla K803.7
Tesla K403.5
Tesla K203.5
Tesla C20752.0
Tesla C2050/C20702.0
NVIDIA Data Center Products
GPUCompute Capability
NVIDIA A1008.0
NVIDIA T47.5
NVIDIA V1007.0
Tesla P1006.0
Tesla P406.1
Tesla P46.1
Tesla M605.2
Tesla M405.2
Tesla K803.7
Tesla K403.5
Tesla K203.5
Tesla K103.0
Quadro Desktop Products
GPUCompute Capability
Quadro RTX 80007.5
Quadro RTX 60007.5
Quadro RTX 50007.5
Quadro RTX 40007.5
Quadro GV1007.0
Quadro GP1006.0
Quadro P60006.1
Quadro P50006.1
Quadro P40006.1
Quadro P22006.1
Quadro P20006.1
Quadro P10006.1
Quadro P6206.1
Quadro P6006.1
Quadro P4006.1
Quadro M6000 24GB5.2
Quadro M60005.2
Quadro K60003.5
Quadro M50005.2
Quadro K52003.5
Quadro K50003.0
Quadro M40005.2
Quadro K42003.0
Quadro K40003.0
Quadro M20005.2
Quadro K22005.0
Quadro K20003.0
Quadro K2000D3.0
Quadro K12005.0
Quadro K6205.0
Quadro K6003.0
Quadro K4203.0
Quadro 4103.0
Quadro Plex 70002.0
Quadro Mobile Products
GPUCompute Capability
RTX 50007.5
RTX 40007.5
RTX 30007.5
T20007.5
T10007.5
P6206.1
P5206.1
Quadro P52006.1
Quadro P42006.1
Quadro P32006.1
Quadro P50006.1
Quadro P40006.1
Quadro P30006.1
Quadro P20006.1
Quadro P10006.1
Quadro P6006.1
Quadro P5006.1
Quadro M5500M5.2
Quadro M22005.2
Quadro M12005.0
Quadro M6205.2
Quadro M5205.0
Quadro K6000M3.0
Quadro K5200M3.0
Quadro K5100M3.0
Quadro M5000M5.0
Quadro K500M3.0
Quadro K4200M3.0
Quadro K4100M3.0
Quadro M4000M5.0
Quadro K3100M3.0
Quadro M3000M5.0
Quadro K2200M3.0
Quadro K2100M3.0
Quadro M2000M5.0
Quadro K1100M3.0
Quadro M1000M5.0
Quadro K620M5.0
Quadro K610M3.5
Quadro M600M5.0
Quadro K510M3.5
Quadro M500M5.0
GeForce and TITAN Products
GPUCompute Capability
GeForce RTX 30908.6
GeForce RTX 30808.6
GeForce RTX 30708.6
NVIDIA TITAN RTX7.5
Geforce RTX 2080 Ti7.5
Geforce RTX 20807.5
Geforce RTX 20707.5
Geforce RTX 20607.5
NVIDIA TITAN V7.0
NVIDIA TITAN Xp6.1
NVIDIA TITAN X6.1
GeForce GTX 1080 Ti6.1
GeForce GTX 10806.1
GeForce GTX 1070 Ti6.1
GeForce GTX 10706.1
GeForce GTX 10606.1
GeForce GTX 10506.1
GeForce GTX TITAN X5.2
GeForce GTX TITAN Z3.5
GeForce GTX TITAN Black3.5
GeForce GTX TITAN3.5
GeForce GTX 980 Ti5.2
GeForce GTX 9805.2
GeForce GTX 9705.2
GeForce GTX 9605.2
GeForce GTX 9505.2
GeForce GTX 780 Ti3.5
GeForce GTX 7803.5
GeForce GTX 7703.0
GeForce GTX 7603.0
GeForce GTX 750 Ti5.0
GeForce GTX 7505.0
GeForce GTX 6903.0
GeForce GTX 6803.0
GeForce GTX 6703.0
GeForce GTX 660 Ti3.0
GeForce GTX 6603.0
GeForce GTX 650 Ti BOOST3.0
GeForce GTX 650 Ti3.0
GeForce GTX 6503.0
GeForce GTX 560 Ti2.1
GeForce GTX 550 Ti2.1
GeForce GTX 4602.1
GeForce GTS 4502.1
GeForce GTS 450*2.1
GeForce GTX 5902.0
GeForce GTX 5802.0
GeForce GTX 5702.0
GeForce GTX 4802.0
GeForce GTX 4702.0
GeForce GTX 4652.0
GeForce GT 7403.0
GeForce GT 7303.5
GeForce GT 730 DDR3,128bit2.1
GeForce GT 7203.5
GeForce GT 705*3.5
GeForce GT 640 (GDDR5)3.5
GeForce GT 640 (GDDR3)2.1
GeForce GT 6302.1
GeForce GT 6202.1
GeForce GT 6102.1
GeForce GT 5202.1
GeForce GT 4402.1
GeForce GT 440*2.1
GeForce GT 4302.1
GeForce GT 430*2.1
GeForce Notebook Products
GPUCompute Capability
Geforce RTX 20807.5
Geforce RTX 20707.5
Geforce RTX 20607.5
GeForce GTX 10806.1
GeForce GTX 10706.1
GeForce GTX 10606.1
GeForce GTX 9805.2
GeForce GTX 980M5.2
GeForce GTX 970M5.2
GeForce GTX 965M5.2
GeForce GTX 960M5.0
GeForce GTX 950M5.0
GeForce 940M5.0
GeForce 930M5.0
GeForce 920M3.5
GeForce 910M5.2
GeForce GTX 880M3.0
GeForce GTX 870M3.0
GeForce GTX 860M3.0/5.0(**)
GeForce GTX 850M5.0
GeForce 840M5.0
GeForce 830M5.0
GeForce 820M2.1
GeForce 800M2.1
GeForce GTX 780M3.0
GeForce GTX 770M3.0
GeForce GTX 765M3.0
GeForce GTX 760M3.0
GeForce GTX 680MX3.0
GeForce GTX 680M3.0
GeForce GTX 675MX3.0
GeForce GTX 675M2.1
GeForce GTX 670MX3.0
GeForce GTX 670M2.1
GeForce GTX 660M3.0
GeForce GT 755M3.0
GeForce GT 750M3.0
GeForce GT 650M3.0
GeForce GT 745M3.0
GeForce GT 645M3.0
GeForce GT 740M3.0
GeForce GT 730M3.0
GeForce GT 640M3.0
GeForce GT 640M LE3.0
GeForce GT 735M3.0
GeForce GT 635M2.1
GeForce GT 730M3.0
GeForce GT 630M2.1
GeForce GT 625M2.1
GeForce GT 720M2.1
GeForce GT 620M2.1
GeForce 710M2.1
GeForce 705M2.1
GeForce 610M2.1
GeForce GTX 580M2.1
GeForce GTX 570M2.1
GeForce GTX 560M2.1
GeForce GT 555M2.1
GeForce GT 550M2.1
GeForce GT 540M2.1
GeForce GT 525M2.1
GeForce GT 520MX2.1
GeForce GT 520M2.1
GeForce GTX 485M2.1
GeForce GTX 470M2.1
GeForce GTX 460M2.1
GeForce GT 445M2.1
GeForce GT 435M2.1
GeForce GT 420M2.1
GeForce GT 415M2.1
GeForce GTX 480M2.0
GeForce 710M2.1
GeForce 410M2.1

CUDA事件:

CUDA中事件本质上是CUDA流中的标记，它与该流内操作流中特定点相关联

可以使用事件来执行以下两个基本任务：

  • 同步流的执行

  • 监控设备的进展

CUDA的API提供了在流中任意点插入事件以及查询事件完成的函数

只有当一个给定CUDA流中先前的所有操作都执行结束后，记录在该流内的事件才会起作用（即完成）

可以理解为:
事件就是在流中插入的一个特殊的操作

创建 & 销毁:

事件的声明&创建&销毁:

cudaEvent_t event;
__host__ cudaError_t CUDARTAPI cudaEventCreate(cudaEvent_t * event);
__host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventDestroy(cudaEvent_t event);

当cudaEventDestroy函数被调用时，如果事件尚未起作用，则调用立即返回，当事件被标记完成时自动释放与该事件相关的资源

记录事件 & 计算运行时间

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

一个事件使用如下函数排队进入CUDA流:

__host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t stream __dv (0))

之后就可以利用此事件检测该事件之前的操作的完成情况:

__host__ cudaError_t CUDARTAPI cudaEventSynchronize(cudaEvent_t event)

此操作会阻塞主机线程

其作用类似于上头的cudaStreamSynchronize, 但其可以等待一个流的中间点(事件插入点)

不阻塞的查询版本:

__host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event)

类似于上头的cudaStreamQuery

而后还有个函数用来计算两个事件完成时间隔的时间:

__host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float * ms, cudaEvent_t start, cudaEvent_t end)

此函数返回事件启动和停止之间的运行时间，以毫秒为单位, 此时事件通常已经完成
startEvent & endEvent不必在同一个CUDA流中

需要注意的是:
使用此函数记录非空流的时间时, 会比实际耗时要长, 因为cudaEventRecord是异步的, 并且不能保证计算的延时整好处在两个事件之间

小栗子:

image-20210216171426544

流同步:

在非默认流中，所有的操作对于主机线程都是非阻塞的，因此会遇到需要在同一个流中同步主机和运算操作的情况

从主机的角度来说，CUDA操作可以分为两大类：

  • 内存相关操作

  • 内核启动

有两种类型的流：

  • 异步流（非空流）

  • 同步流（空流/默认流)

    其又可以细分为以下两种:

    • 阻塞流

    • 非阻塞流

虽然非空流在主机上是非阻塞的，但是非空流内的操作可以被空流中的操作所阻塞

  • 如果一个非空流是阻塞流，则空流可以阻塞该非空流中的操作
  • 如果一个非空流是非阻塞流，则它不会阻塞空流中的操作

在下面的部分中，将介绍如何使用阻塞流和非阻塞流。

阻塞流和非阻塞流:

前头使用的cudsStreamCreate()创建的就是阻塞流

阻塞流和空流的相互依赖关系(同步关系)如下:

当操作被发布到空流中，在该操作被执行之前，CUDA上下文会等待所有先前的操作发布到所有的阻塞流中。此外，任何发布到阻塞流中的操作，会被挂起等待，直到空流中先前的操作执行结束才开始执行

这段有点绕, 但看下头的例子:

image-20210216191611852

这段代码的结果是, 直到核函数kernel_1执行结束，kernel_2才会在GPU上开始执行，kernel_2执行结束后，kernel_3才开始执行

创建非阻塞流:

__host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t * pStream, unsigned int flags)

使用这个函数创建:

flags参数用于指定是否创建非阻塞流:

image-20210216192234584

指定cudaStreamNonBlocking使得非空流对于空流的阻塞行为失效

如果前头的例子中使用非阻塞流, 则所有的核函数执行都不会被阻塞, 都不用等待其他核函数执行结束

隐式同步:

这里就是介绍几个会发生隐式同步的地方

隐式同步在CUDA编程中特别吸引编程人员的注意，因为带有隐式同步行为的运行时函数可能会导致不必要的阻塞，这种阻塞通常发生在设备层面

例如：

  • 锁页主机内存分配

  • 设备内存分配

  • 设备内存初始化

  • 同一设备上两个地址之间的内存复制

  • 一级缓存/共享内存配置的修改

显示同步:

之前介绍了几种 主机-设备的显示同步:

cudaDeviceSynchronize，cudaStreamSynchronize & cudaEventSynchronize

这里再介绍几个多stream同步的方法:

  • 同步设备

  • 同步流

  • 同步流中的事件

  • 使用事件跨流同步

阻塞主机线程直到设备完成所有先前的任务:

__host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceSynchronize(void)

阻塞主机线程直到指定的流中所有任务完成

还有一个查询函数

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

阻塞主机线程直到指定的事件完成

还有个查询函数

__host__ cudaError_t CUDARTAPI cudaEventSynchronize(cudaEvent_t event);
__host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event);

此外, 还有个比较灵活的方法:

__host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags __dv (0))

该函数会指定该stream等待特定的event, 仅有该event触发之后才会启动流
该event可以关联到相同或者不同的stream

所以要求是调用这个函数前, 不能向指定的流添加任何任务

image-20210216194856854

可配置事件:

这里具体还不是很清楚其他flags 的意义

CUDA运行时提供了一种方式来定制事件的行为和性能:

__host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t * event, unsigned int flags)

image-20210216194939501

书里翻译的极其奇怪, 这里使用官方CUDA源文:

Creates an event object for the current device with the specified flags. Valid flags include:

  • default就是创建一个普通的事件

  • cudaEventBlockingSync创建一个会阻塞host 的event
    即指定了所创建的event应该使用cudaEventSynchronize()等待事件完成, 而不是异步进行

  • cudaEventDisableTiming创建一个只用来同步, 而不用来计算时间的event

    这样就能减少记录时间戳的消耗, 提高cuudaStreamWaitEvent和cudaEventQuery的调用性能

  • cudaEventInterprocess指定event可以被用来作为inter-process event
    这玩意暂时用不到

6.2 并发内核执行:

这一节就是将上头讲到的几个例子实践一波:

包括以下几个方面：

  • 使用深度优先或广度优先方法的调度工作

  • 调整硬件工作队列

  • 在Kepler设备和Fermi设备上避免虚假的依赖关系

  • 检查默认流的阻塞行为

  • 在非默认流之间添加依赖关系

  • 检查资源使用是如何影响并发的

非空流中的并发内核

在本节中，将使用NVIDIA的可视化性能分析器（nvvp）可视化并发核函数执行

为了能更好的体现出差异性, 核函数在GPU上驻留的时间要足够长, 所以本例中使用了多个完全相同的核函数:

image-20210216202637207

创建一组非空流, 并为每个流添加任务:

image-20210216203935469

image-20210216204201710

并加入到默认流中:

image-20210216204232248

image-20210216204250464

输出结果:

在Tesla K40上运行

image-20210216204310686

image-20210216204332322

Fermi GPU上的虚假依赖关系:

由于基本不可能使用到Fermi架构的GPU了, 所以这节可略

但是为了理解虚假依赖关系, 还是看看较好

image-20210216204543783

image-20210216204643547

这个时序图和上一节讲的相同, 仅有最后一个任务和另一个流的第一个任务能并行

避免虚假依赖关系:

image-20210216204757612

image-20210216204804964

image-20210216204812460

image-20210216204809370

使用OpenMP的调度操作:

本小节使用OpenMP创建多个主机线程进行调度

这里就是将上一个小节中的循环调度操作改用OpenMP进行多线程并行

image-20210216205817067

image-20210216205830432

可以看到这种简单的并行并没有带来明显的性能差异

仅有当主机&设备都有计算任务时, 主机上的多线程才能更加有效

用环境变量调整流行为

之前讲到Kepler设备支持的硬件工作队列的最大数量是32, 然而，默认情况下并发硬件连接的数量被限制为8
由于每个连接都需要额外的内存和资源，所以设置默认的限制为8，减少了不需要全部32个工作队列的应用程序的资源消耗

可以使用CUDA_DEVICE_MAX_CONNECTIONS环境变量来调整并行硬件连接的数量，对于Kepler设备而言，其上限是32

有几种设置该变量的方法:

  1. Linux中使用Bash & Bourne Shell:

    image-20210216210219740

  2. C-Shell:

    image-20210216210236409

  3. 直接在主机程序中进行设定:

    image-20210216210251938

修改上头的程序以呈现出流数量大于工作队列数量的情况:

image-20210216210431576

image-20210216210406727

图6-10展示了8个流，但是只有4路并发。因为现在只有4个设备连接，两个流共享一个队列

GPU资源的并发限制:

有限的内核资源可以抑制应用程序中可能出现的内核并发的数量

在之前的例子中，由于避免出现资源导致的并发限制, 所以启动内核时只有一个线程
在本例中, 将提高核函数的运行时配置, 以测试资源导致的并发限制

image-20210216210643202

image-20210216210649877

image-20210216210700870

如图6-12所示，图中只实现了8路并发，即使CUDA设备连接的数量被设置为32。因为GPU无法分配足够的资源来执行所有符合条件的内核，所以并发性是有限的

所以可知, 即使工作队列的上限是32, 但是由于硬件设备资源的原因, 实际的最大并发流个数还是需要实际测试的

默认流的阻塞行为:

本小节中说明了默认流在非空流中是如何阻塞操作的

将上头例子中的主机程序修改如下:

image-20210216210911262

因为第三个内核在默认流中被启动，所以在非空流上所有之后的操作都会被阻塞，直到默认流中的操作完成

image-20210216210935939

创建流间依赖关系:

在理想情况下，流之间不应该有非计划之内的依赖关系（即虚假的依赖关系）
但在复杂的应用程序中，不同的流之间的工作很可能需要相互配合

本例就是利用上头的cudaStreamWaitEvent创建流之间的依赖关系

假设想让一个流中的工作在其他所有流中的工作都完成后才开始执行

image-20210216211259843

而后将这些事件插入到其他流的末尾, 用以标志流的结束, 并使最后一个流等待这些事件的完成:

image-20210216211421295

image-20210216211427623

6.3 重叠内核执行 & 数据传输:

在前一节中，已经介绍了如何在多个流中并发执行多个内核。在本节中，将学习如何并发执行内核和数据传输

重叠内核和数据传输表现出不同的行为，并且需要考虑一些与并发内核执行相比不同的因素

Fermi GPU和Kepler GPU有两个复制引擎队列：

  • 一个用于将数据传输到设备
  • 另一个用于从设备中将数据提取出来

所以最多可以重叠两个不同方向的数据传输, 否则, 所有的数据传输都将是串行的

在应用程序中，还需要检验数据传输和内核执行之间的关系，从而可以区分以下两种情况：

  • 如果一个内核使用数据A，那么对A进行数据传输必须要安排在内核启动前，且必须位于相同的流中

    这种情况下，要实现数据传输和内核执行之间的重叠会更复杂，因为内核依赖数据作为输入

  • 如果一个内核完全不使用数据A，那么内核执行和数据传输可以位于不同的流中

    这种情况下，实现内核和数据传输的并发执行是很容易的：将它们放置在不同的流中，这就已经向运行时表示了并发地执行它们是安全的

使用深度优先调度重叠:

本例中将使用上头熟悉的向量加法核函数:

image-20210216212322289

其中增大了n_repeat从而以延长内核的执行时间

实现向量加法的CUDA程序，其基本结构包含3个主要步骤：

  • 将两个输入向量从主机复制到设备中

  • 执行向量加法运算

  • 将单一的输出向量从设备返回主机中

之前执行的操作是将全部的数据一次性拷贝到设备内存中, 而后在执行计算, 最后再将结果一次性拷贝回来

而现在将向量加法划分为多个子任务, 这样就能实现不同子任务的数据拷贝和数据计算的重叠

要重叠数据传输和内核执行，必须使用异步复制函数

image-20210216213628424

接下来需要将n个子任务分配到n个流中:

image-20210216213700752

现在，可以使用一个循环来为几个流同时调度iElem个元素的通信和计算

image-20210216213745555

通过将数据传输和该数据上的计算放在同一个流中，输入向量、内核计算以及输出向量之间的依赖关系可以被保持

并且由于多个子任务互不影响, 所以可以分配到多个流中并行执行

为了进行对比，此例还使用了一个阻塞实现来计算基准性能：

image-20210216213940583

image-20210216214010981

图6-15显示了Tesla K40设备的时间轴。图中使用了8个硬件工作队列和4个CUDA流来重叠内核执行和数据传输
相对于阻塞的默认流执行，该流执行实现了近40%的性能提升

图6-15显示了以下3种重叠：

  • 不同流中内核的互相重叠

  • 内核与其他流中的数据传输重叠

  • 在不同流以及不同方向上的数据传输互相重叠

图6-15还呈现了以下两种阻塞行为：

  • 内核被同一流中先前的数据传输所阻塞

  • 从主机到设备的数据传输被同一方向上先前的数据传输所阻塞

网格管理单元:

Kepler引入了一个新的网格管理和调度控制系统，即网格管理单元（GMU）

GMU可以暂停新网格的调度，使网格排队等待且暂停网格直到它们准备好执行，这样就使运行时变得非常灵活强大，动态并行就是一个很好的例子

在Fermi设备上，网格直接从流队列被传到CUDA工作分配器（CUDA Work Distributor，CWD）中。在Kepler设备上，网格被发送到GMU上，GMU对在GPU上执行的网格进行管理和优先级排序

GMU创建了多个硬件工作队列，从而减少或消除了虚假的依赖关系。通过GMU，流可以作为单独的工作流水线。即使GMU被限制只能创建一个单一的硬件工作队列，根据以上测试结果证实，通过GMU进行的网格依赖性分析也可以帮助消除虚假的依赖关系

使用广度优先调度重叠:

修改核函数如下以实现广度优先的效果

image-20210216214920615

image-20210216215001374

图6-17显示了在K40设备上只使用一个硬件工作队列时的时间轴

与深度优先的方法相比它没有明显的差异，因为Kepler的双向调度机制有助于消除虚假的依赖关系
但如果在Fermi设备上运行相同的测试，在整体性能方面会发现，使用广度优先方法不如使用深度优先方法。由主机到设备复制队列上的争用导致的虚假依赖关系，在主机到设备间的传输完成前，将阻止所有的内核启动

所以结论是:

  • 对于Kepler架构的GPU, 通常无需关注其工作调度顺序
  • 对于Fermi设备, 则需要注意这些问题

6.4 重叠GPU和CPU的执行:

相对而言，实现GPU和CPU执行重叠是比较简单的，因为所有的内核启动在默认情况下都是异步的

只需简单地启动内核，并且立即在主机线程上实现有效操作，就能自动重叠GPU和CPU执行

本节的示例主要包括两个部分：

  • 内核被调度到默认流中

  • 等待GPU内核时执行主机计算

使用以下简单的内核实现一个向量与标量的加法：

image-20210216215217724

本例中使用了3个CUDA操作（两个复制和一个内核启动）。记录一个停止事件，以标记所有CUDA操作的完成。

image-20210216215433964

由于上头使用的全是异步操作, 所以在这之后控制权立即返回到主机, 此时即可进行计算操作:
这里做的工作就是计数器不断++

image-20210216215557850

以下是在Tesla K40上使用nvprof的输出。在等待GPU操作完成时，主机线程执行14606次迭代。

image-20210216215630640

6.5 流回调

回调这个玩意应该不陌生, 所以能很快的猜到:

相当于往CUDA流中塞入了自建函数而非CUDA API
即自建函数与CUDA API一样能在CUDA流中排队执行

回调功能十分强大，因为它们是第一个GPU操作的例子，此操作与之前所学的都相反, 是GPU在主机上创建任务

流回调函数是由应用程序提供的一个主机函数，其有特殊的格式:
(由于C++的函数指针都有带参数的, 所以必须格式相同)

void CUDART_CB my_callback(cudaStream_t stream, cudaError_t status, void *data) {
    //Do something
}

并在流中使用以下的API函数注册:

__host__ cudaError_t CUDARTAPI cudaStreamAddCallback(cudaStream_t stream, cudaStreamCallback_t callback, void * userData, unsigned int flags)

其中:

  • callback:
    主机端指定的回调函数
  • userData
    就是传递给回调函数的参数
  • flags:

回调函数有俩限制:

  • 从回调函数中不可以调用CUDA的API函数
  • 在回调函数中不可以执行同步

回调函数的俩特点:

  • 每使用cudaStreamAddCallback一次，只执行一次回调
  • 回调函数执行时, 阻塞队列中排在其后面的工作，直到回调函数完成

小栗子:

image-20210216221037514

总结:

流的概念是CUDA编程模型的一个基本组成部分。允许高级CUDA操作在独立的流中排队执行，CUDA流可以实现粗粒度并发。因为CUDA支持异步操作和大多数版本的运行时函数，所以它可以在多个CUDA流之间调度计算和通信。

从概念上讲，如果CUDA操作之间存在依赖关系，则它们必须在同一个流中被调度。例如，为了确保应用程序的准确无误，内核必须在同一个流中被调度，并在它使用的任何数据传输后进行。另外，没有依赖关系的操作可以在任意的流中被调度。在CUDA中，通常可以使用3种不同类型的重叠方案来隐藏计算或通信延迟：

  • 在设备上重叠多个并发的内核

  • 重叠带有传入或传出设备数据传输的CUDA内核

  • 重叠CPU执行和GPU执行

为了充分利用设备，并确保最大的并发性，还需要注意以下问题：

  • 平衡内核资源需求和并发资源需求。在设备上一次启动过多的计算任务，可能会导致内核串行，这会使得硬件资源的工作块变得可用。但是，也需要确保设备没有被充分利用，一直有工作在排队等待执行。

  • 如果可能的话，避免使用默认流执行异步操作。放置在默认流中的操作可能会阻塞其他非默认CUDA流的进展。

  • 在Fermi设备上，从深度优先和广度优先两方面考虑主机的调度。这个选择可以通过消除共享硬件工作队列上的虚假依赖关系，显著影响其性能。

  • 要注意隐式同步的函数，并且充分利用它们和异步函数来避免性能的降低。

此外，本章还介绍了CUDA可视化性能分析器（nvvp）在可视化GPU执行中的作用。nvvp允许确认操作重叠的条件，并且易于多个流行为的可视化。

Janus_V
关注
  • 7
    点赞
  • 踩
  • 28
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
专栏目录
Cuda异步计算并行编程原理和存储管理
tugouxp的专栏
12-17 1599
优化时遵守的一些原则：1.Grid一定要给足block.2.Block内一定要给足thread,目的是提高并发WARP的数目，隐藏延迟.3.Block内线程的数目一定是warpsize的整数倍.每个warp的执行上下文（execution context，如程序计数器 和 寄存器等）在warp的整个生命周期内都被保存在片上内存（on-chip memory）。因此从一个执行上下文切换到另一个执行上下文是无开销的。
CUDA流任务并行
xysjj的博客
12-19 3199
CUDA流表示一个GPU操作队列，并且该队列中的操作将以指定的顺序执行。可以将每个流视为GPU的一个任务，并且这些任务可以并行执行，即相同流顺序执行，不同流并行执行；不同流并行执行时不同流所要执行的任务要没有依赖关系；当不手动创建流时，cuda将会默认一个流操作。 在硬件选择上，这里有一个概念，支持设备重叠功能，支持设备重叠功能的 GPU 能够在执行一个 CUDA C...
CUDA流：利用并行执行提高性能
Algabeno的博客
02-22 3738
CUDA流说明
cuda笔记-流的使用（定义、创建、消耗、同步）
IT1995的博客
12-29 1507
定义流： cudaStream_t stream; 创建流： cudaStreamCreate(cudaStream_t *s) 销毁流： cudaStreamDestroy(cudaStream_t s) 流的同步：显示同步 //同步单个流，等待该命令都完成 cudaStreamSynchronize(cudaStream_t stream); //同步所有流，等待整个设备上的流都完成 cudaDeviceSynchronize(); //通过某个事件：等待某个事件结束后执行该流
我与计算机视觉-[CUDA]-[CPU多线程下CUDA的多流]
weixin_43409627的博客
07-29 2246
首先问题出在cpu下的多线程，当你想要在多个线程下调用同一个cuda核函数的时候，你会发现效率很低，那么经过验证，的确，不管你有多少个线程，cuda总是将线程中的核函数放入默认流中进行队列方式的处理，相当于单线程，但是这个问题在cuda7后已经得到了解决，下面对这个问题进行一些测试。 这里使用了以下链接中的内容： https://www.cnblogs.com/wujianming-110117/p/14091897.html https://developer.nvidia.com/blog/gpu
GPU编程 CUDA C++ 使用多个非默认CUDA流并发重叠核函数执行和数据传输
行子的博客
04-03 440
1. 从主机到GPU的数据传输量（d_x，d_y）是从GPU到主机的数据传输量（d_z）的两倍。使用约20个非默认CUDA流并发重叠核函数执行和数据传输，有着最大的性能，得到了约2倍的加速比。继续增加并发CUDA流数量n后的极限时间是：(n+2)/(3n) 约为 1/3。使用四个非默认CUDA流并发重叠：（时间是：6/12）2. 当CUDA流的数量增加，会带来额外的开销。使用一个默认CUDA流：（时间是：1）D：Device设备（GPU）H2D：主机复制数据到GPU。D2H：GPU复制数据到主机。
cuda-tutorial:CUDA编程简短教程
05-18
总之，CUDA编程是解锁GPU计算潜力的关键，通过学习和实践，开发者能够编写出高效、并行的代码，充分利用现代GPU的计算能力。这个简短教程是探索这一领域的良好起点，后续需要结合更多资源和实践项目进一步提升技能。
CUDA编程学习（六份pdf）
05-03
CUDA编程学习3.pdf可能介绍了如何定义和控制线程块和网格，以及如何进行线程间的同步，以确保正确执行并发任务。 第四部分：流与设备管理 CUDA流允许异步执行，提高资源利用率。设备管理则涉及数据在GPU和CPU之间的...
CUDA 编程：基础与实践_樊哲勇1
08-03
接着，第6章至第8章详细阐述了CUDA的内存系统，包括全局内存、共享内存、常量内存和纹理内存的使用。第9章介绍了原子操作，这是在并行环境中实现同步和数据一致性的重要工具。第10章讲解了线程束内的基本函数，如...
CUDA编程指南5.0
08-29
第一章导论 1 1.1 从图形处理到通用并行计算 1 1.2 CUDATM：一种通用并行计算架构 3 1.3 一种可扩展的编程模型 3 ...iv CUDA编程指南5.0中文版 5.4.1 算术指令 101 5.4.2 控制流指令 104 5.4.3 同步指令 ...
cuda编程002—流
最新发布
Little__Black的博客
08-23 300
cuda流整体笔记和代码。
VS2017 CUDA编程学习12：CUDA流
DU_YULIN的博客
11-12 3049
文章目录前言1. CUDA流的理解2. C++ 实现CUDA流3. 执行结果总结学习资料 VS2017 CUDA编程学习1：CUDA编程两变量加法运算 VS2017 CUDA编程学习2：在GPU上执行线程 VS2017 CUDA编程学习3：CUDA获取设备上属性信息 VS2017 CUDA编程学习4：CUDA并行处理初探 - 向量加法实现 VS2017 CUDA编程学习5：CUDA并行执行-线程 VS2017 CUDA编程学习6: GPU存储器架构 VS2017 CUDA编程学习7：线程同步-共享内存 VS
CUDA编程入门极简教程
热门推荐
小白将
03-18 9万+
码字不易，欢迎给个赞！ 欢迎交流与转载，文章会同步发布在公众号：机器学习算法全栈工程师(Jeemy110) 目录 目录 前言 CUDA编程模型基础 向量加法实例 矩阵乘法实例 小结 参考资料 前言 2006年，NVIDIA公司发布了CUDA，CUDA是建立在NVIDIA的CPUs上的一个通用并行计算平台和编程模型，基于CUDA编程可以利用GPUs的并行计算引擎...
Cuda Streams的概述（四）-- 同步
大脸猫脸大主持人猫脸大的博客
05-17 4801
同步 同步的APIs 同步所有的事情 //阻塞host端，直到所有的CUDA调用完成。 cudaDeviceSynchronize(); 同步主机端特定的流 //阻塞host端，直到流里的CUDA调用完成。 cudaStreamSynchronize(stream); 通过主机端或设备用events CUDA Events 当操作在流中发生时，提供一个信号机制。这个对profiling和同步是很有用的。 Event有一个boolean状态： 发生 未发生 注意：默认状态=发生 管理Event //
CUDA的kernel并行
feng__shuai的博客
01-11 3629
流并行
CUDA学习之流和并发--part1
weixin_42764932的博客
04-07 1167
文章目录6.1 流和事件概述6.1.1 CUDA流6.1.2 流调度1. 虚假的依赖关系2. Hyper-Q技术6.1.3 流的优先级CUDA事件1. 创建和销毁2. 记录事件和计算运行时间6.1.5 流同步1. 阻塞流和非阻塞流2. 隐式同步3. 显式同步4. 可配置事件6.2 并发内核执行6.2.1 非空流中的并发内核6.2.2 Fermi GPU 上的虚假依赖关系6.2.3 使用OpenMP的调度操作6.2.4 用环境变量调整流行为6.2.5 GPU资源的并发限制6.2.6 默认流的阻塞行为6.2.7
CUDA学习笔记（3）- 流并行和线程同步
不会飞的纸飞机的博客
09-17 3061
文章目录1. 流并行2. 线程同步 1. 流并行 线程流中可以有多个线程块，线程块中可以有多个线程。线程块和线程流只能处理单个函数，线程流可以处理多个函数和同一个函数的不同参数。 cudaStreamCreate(cudaStream_t *pStream) 创建一个线程流。 cudaStreamDestroy(cudaStream_t stream) 销毁线程流。 下面是关于流并行的简...
CUDA异步并发 （二）
weixin_33725126的博客
10-15 580
感觉自己对CUDA的异步并发还是无法深入理解，在大神建议下看了CUDA C Programming Guide ，感觉对自己的帮助很大。 下面分享一下看了之后的体会，附上原文： Concurrent host execution is facilitated through asynchronous library function...
GPU 高性能编程 CUDA : 流
xxiaozr的博客
05-26 6816
对于大规模数据的并行运算，GPU上的执行性能远高于 CPU 上的性能，除此之外， NVIDAI 图像处理器还支持另一种类型的并行性，类似于 CPU 的多线程应用程序中的任务并行性。任务并行性是指并行执行两个或者多个任务，而不是在大量数据上执行同一个任务。页锁定主页内存之前都是使用 cudaMalloc() 在 GPU 上分配内存，以及通过标准 C 库函数 malloc() 在主机上分配内存除此之外...
用CUDA代码实现cv::seamlessClone
08-05
### 回答1： 实现cv::seamlessClone可以使用OpenCV库中提供的CUDA函数进行实现。 以下是一个简单的示例代码： ``` #include <opencv2/opencv.hpp> #include <opencv2/cudaimgproc.hpp> #include <opencv2/cudaarithm.hpp> int main(int argc, char** argv) { cv::Mat src = cv::imread("src.jpg"); cv::Mat dst = cv::imread("dst.jpg"); cv::Mat mask = cv::imread("mask.jpg", 0); cv::cuda::GpuMat src_gpu, dst_gpu, mask_gpu, result_gpu; src_gpu.upload(src); dst_gpu.upload(dst); mask_gpu.upload(mask); cv::cuda::seamlessClone(src_gpu, dst_gpu, mask_gpu, cv::Point(dst.cols / 2, dst.rows / 2), result_gpu, cv::cuda::NORMAL_CLONE); cv::Mat result_cpu; result_gpu.download(result_cpu); cv::imshow("result", result_cpu); cv::waitKey(0); return 0; } ``` 在此代码中，我们首先加载了原始图像、目标图像和掩码图像，然后将它们上传到GPU。接下来，我们调用`cv::cuda::seamlessClone`函数，并将结果下载到CPU上的矩阵中。最后，我们使用`cv::imshow`函数显示结果。 ### 回答2： 使用CUDA代码实现cv::seamlessClone需要以下步骤： 1. 首先，将输入图像和目标图像从主机内存复制到CUDA设备内存中。可以使用cudaMemcpy函数进行内存拷贝。 2. 在CUDA设备上创建一个输出图像的内存空间，并使用cudaMalloc函数为其分配内存。 3. 将输入图像和目标图像的像素数据分别传送到CUDA设备内存中。可以使用cudaMemcpy2D函数将二维图像数据传送到设备。 4. 在CUDA设备上创建一个内核函数，用来计算图像中的每个像素点的融合颜色。该函数可以根据融合算法的不同，使用不同的插值方法来计算像素点的新颜色。 5. 调用内核函数，对每个像素点进行并行计算，计算结果存储在输出图像内存中。 6. 最后，将输出图像的像素数据从设备内存复制到主机内存中。可以使用cudaMemcpy2D函数将二维图像数据从设备复制到主机内存。 7. 在主机上，创建一个新的cv::Mat对象，并将复制的像素数据填充到该对象中。最后，在主机上释放设备内存。 需要注意的是，实现CUDA版本的cv::seamlessClone可能需要一些图像处理和计算机视觉的知识，以及对CUDA编程模型的理解。同时，需要具备使用CUDA编程环境和库函数的能力。 ### 回答3： cv::seamlessClone函数是OpenCV中用于图像无缝融合的函数。要使用CUDA代码实现类似的功能，可以参考以下步骤： 1. 从输入图像和目标图像中读取数据，并将其分配到CUDA设备的全局内存中。 2. 创建一个与输入图像和目标图像大小相同的空白图像作为输出图像，并将其分配到CUDA设备的全局内存中。 3. 在CUDA设备上为输入图像、目标图像和输出图像分配相应的内存空间。 4. 使用CUDA核函数对输入图像和目标图像进行处理，计算图像的梯度（通过Sobel算子或其他方法），并将结果存储在CUDA设备内存中。 5. 使用CUDA核函数对输出图像进行处理，将输入图像和目标图像的梯度信息以及融合参数（比如像素权重）进行计算，并在输出图像中生成无缝融合的效果。 6. 将输出图像从CUDA设备的内存复制到主机内存，以便进一步处理或保存。 7. 释放CUDA设备内存中的图像数据和其他资源。 通过以上步骤，就可以用CUDA代码实现类似于cv::seamlessClone函数的功能，实现图像的无缝融合。但是具体的实现需要根据具体的需求和使用情况来进行一些调整和优化，以提高算法的效率和准确性。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 1
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

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

余额充值