本章内容：
-
理解流和事件的本质
-
理解网格级并发
-
重叠内核执行和数据传输
-
重叠CPU和GPU执行
-
理解同步机制
-
避免不必要的同步
-
调整流的优先级
-
注册设备回调函数
-
通过NVIDIA可视化性能分析器显示应用程序执行的时间轴
一般来说，在CUDA C编程中有两个级别的并发：
-
内核级并发
-
网格级并发
到目前为止，你的关注点可能仅限于内核级的并发，在此级别的并发中，单一的任务或内核被GPU的多个线程并行执行。前面几章已经介绍了提升内核性能的几种方法，它们分别是从编程模型、执行模型和内存模型的角度进行介绍的。想必你已经了解了一些通过命令行性能分析器来研究和分析内核行为的方法。
本章将研究网格级的并发。在网格级并发中，多个内核在同一设备上同时执行，这往往会让设备利用率更好。在本章中，你将学习到如何使用CUDA流实现网格级的并发。还将使用CUDA的可视化性能分析器nvvp将内核并发执行可视化
6.1 流和事件概述:
CUDA流是一系列异步的CUDA操作，这些操作按照主机代码确定的顺序在设备上执行
CUDA流更像是一个由操作组成的队列, 流中的操作按照指定的顺序执行, 与主机保持异步状态, 并可查询队列状态
流中的操作是按照严格的顺序执行的, 但是不同的流中的操作执行顺序则不受限制
所以使用多个流同时启动多个内核, 可以实现网格级的并发, 并且能隐藏这些操作的延时
在许多情况下，执行内核比传输数据耗时更多。在这些情况下，可以完全隐藏CPU和GPU之间的通信延迟。通过将内核执行和数据传输调度到不同的流中，这些操作可以重叠，程序的总运行时间将被缩短。流在CUDA的API调用粒度上可实现流水线或双缓冲技术
CUDA流:
所有的CUDA操作（包括内核和数据传输）都在一个流中显式或隐式地运行
流分为两种类型：
-
隐式声明的流（空流）
-
显式声明的流（非空流）
如果没有显式地指定一个流，那么内核启动和数据传输将默认使用空流
所以之前使用的都是空流或默认流
另一方面, 非空流可以被显式地创建和管理
并且, 如果想要重叠不同的CUDA操作，必须使用非空流
基于流的异步的内核启动和数据传输支持以下类型的粗粒度并发：
-
重叠主机计算和设备计算
-
重叠主机计算和主机与设备间的数据传输
-
重叠主机与设备间的数据传输和设备计算
-
并发设备计算
参考之前的默认流代码:
- 两个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)
这里需要额外注意的是, 使用异步数据传输时, 必须使用固定(非分页)的主机内存
流的使用
在核函数的启动上, 也与使用默认流不同:
需要提供一个额外的流标识符作为第四个参数
流的销毁
使用完流之后, 也可以手动销毁:
当流中的工作尚未完成时, 销毁失败, cudaStreamDestroy函数将立刻返回
只有当流内的操作全部完成时, 才会正常销毁
流的查询:
因为所有的CUDA流操作都是异步的，所以CUDA的API提供了两个函数来检查流中所有操作是否都已经完成
-
cudaStreamSynchronize强制阻塞主机，直到在给定流中所有的操作都完成了
-
cuda-StreamQuery会检查流中所有操作是否都已经完成，但在它们完成前不会阻塞主机:
当所有操作都完成时cudaStreamQuery函数会返回cudaSuccess
当一个或多个操作仍在执行或等待执行时返回cudaErrorNotReady
一个简单的栗子:
以下展现了在多个流中调度CUDA操作的栗子
在图6-1中，数据传输操作虽然分布在不同的流中，但是并没有并发执行。这是由一个共享资源导致的：PCIe总线。虽然从编程模型的角度来看这些操作是独立的，但是因为它们共享一个相同的硬件资源，所以它们的执行必须是串行的。具有双工PCIe总线的设备可以重叠两个数据传输，但它们必须在不同的流中以及不同的方向上。在图6-1中可以观察到，在一个流中从主机到设备的数据传输与另一个流中从设备到主机的数据传输是重叠的
并发内核的最大数量是依赖设备而确定的
Fermi设备支持16路并发，Kepler设备支持32路并发。设备上可用的计算资源进一步限制了并发内核的数量，如共享内存和寄存器。在本章后面的例子中将会探索这些局限性
流调度:
从逻辑上讲, 所有的流都可以并发执行
但是受限于物理硬件, 流需要合理的调度才能正确的执行
虚假的依赖关系:
这里需要介绍CUDA运行时的一个特点:
工作队列中，一个被阻塞的操作会将队列中该操作后面的所有操作都阻塞，即使它们属于不同的流
这种不同流的任务之间的依赖关系就是虚假依赖关系
如图所示:
在流的工作队列中, 后续工作需要等待前头的工作完成之后才能进行, 这种就是依赖关系(有点类似工序图)
所以如果按照图中输入的三个流, 仅有画圈的部分能够被并行, 其余的都是串行执行
Hyper-Q技术:
Kepler架构中引入的Hyper-Q技术
其使用多个硬件工作队列，从而减少了虚假的依赖关系
Hyper-Q技术通过在主机和设备之间维持多个硬件管理上的连接，允许多个CPU线程或进程在单一GPU上同时启动工作
被Fermi架构中虚假依赖关系限制的应用程序，在不改变任何现有代码的情况下可以看到显著的性能提升
但需要注意的是, 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
| GPU | Compute Capability |
|---|---|
| Tesla K80 | 3.7 |
| Tesla K40 | 3.5 |
| Tesla K20 | 3.5 |
| Tesla C2075 | 2.0 |
| Tesla C2050/C2070 | 2.0 |
NVIDIA Data Center Products
| GPU | Compute Capability |
|---|---|
| NVIDIA A100 | 8.0 |
| NVIDIA T4 | 7.5 |
| NVIDIA V100 | 7.0 |
| Tesla P100 | 6.0 |
| Tesla P40 | 6.1 |
| Tesla P4 | 6.1 |
| Tesla M60 | 5.2 |
| Tesla M40 | 5.2 |
| Tesla K80 | 3.7 |
| Tesla K40 | 3.5 |
| Tesla K20 | 3.5 |
| Tesla K10 | 3.0 |
Quadro Desktop Products
| GPU | Compute Capability |
|---|---|
| Quadro RTX 8000 | 7.5 |
| Quadro RTX 6000 | 7.5 |
| Quadro RTX 5000 | 7.5 |
| Quadro RTX 4000 | 7.5 |
| Quadro GV100 | 7.0 |
| Quadro GP100 | 6.0 |
| Quadro P6000 | 6.1 |
| Quadro P5000 | 6.1 |
| Quadro P4000 | 6.1 |
| Quadro P2200 | 6.1 |
| Quadro P2000 | 6.1 |
| Quadro P1000 | 6.1 |
| Quadro P620 | 6.1 |
| Quadro P600 | 6.1 |
| Quadro P400 | 6.1 |
| Quadro M6000 24GB | 5.2 |
| Quadro M6000 | 5.2 |
| Quadro K6000 | 3.5 |
| Quadro M5000 | 5.2 |
| Quadro K5200 | 3.5 |
| Quadro K5000 | 3.0 |
| Quadro M4000 | 5.2 |
| Quadro K4200 | 3.0 |
| Quadro K4000 | 3.0 |
| Quadro M2000 | 5.2 |
| Quadro K2200 | 5.0 |
| Quadro K2000 | 3.0 |
| Quadro K2000D | 3.0 |
| Quadro K1200 | 5.0 |
| Quadro K620 | 5.0 |
| Quadro K600 | 3.0 |
| Quadro K420 | 3.0 |
| Quadro 410 | 3.0 |
| Quadro Plex 7000 | 2.0 |
Quadro Mobile Products
| GPU | Compute Capability |
|---|---|
| RTX 5000 | 7.5 |
| RTX 4000 | 7.5 |
| RTX 3000 | 7.5 |
| T2000 | 7.5 |
| T1000 | 7.5 |
| P620 | 6.1 |
| P520 | 6.1 |
| Quadro P5200 | 6.1 |
| Quadro P4200 | 6.1 |
| Quadro P3200 | 6.1 |
| Quadro P5000 | 6.1 |
| Quadro P4000 | 6.1 |
| Quadro P3000 | 6.1 |
| Quadro P2000 | 6.1 |
| Quadro P1000 | 6.1 |
| Quadro P600 | 6.1 |
| Quadro P500 | 6.1 |
| Quadro M5500M | 5.2 |
| Quadro M2200 | 5.2 |
| Quadro M1200 | 5.0 |
| Quadro M620 | 5.2 |
| Quadro M520 | 5.0 |
| Quadro K6000M | 3.0 |
| Quadro K5200M | 3.0 |
| Quadro K5100M | 3.0 |
| Quadro M5000M | 5.0 |
| Quadro K500M | 3.0 |
| Quadro K4200M | 3.0 |
| Quadro K4100M | 3.0 |
| Quadro M4000M | 5.0 |
| Quadro K3100M | 3.0 |
| Quadro M3000M | 5.0 |
| Quadro K2200M | 3.0 |
| Quadro K2100M | 3.0 |
| Quadro M2000M | 5.0 |
| Quadro K1100M | 3.0 |
| Quadro M1000M | 5.0 |
| Quadro K620M | 5.0 |
| Quadro K610M | 3.5 |
| Quadro M600M | 5.0 |
| Quadro K510M | 3.5 |
| Quadro M500M | 5.0 |
GeForce and TITAN Products
GeForce Notebook Products
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是异步的, 并且不能保证计算的延时整好处在两个事件之间
小栗子:
流同步:
在非默认流中，所有的操作对于主机线程都是非阻塞的，因此会遇到需要在同一个流中同步主机和运算操作的情况
从主机的角度来说，CUDA操作可以分为两大类：
-
内存相关操作
-
内核启动
有两种类型的流：
-
异步流（非空流）
-
同步流（空流/默认流)
其又可以细分为以下两种:
-
阻塞流
-
非阻塞流
-
虽然非空流在主机上是非阻塞的，但是非空流内的操作可以被空流中的操作所阻塞
- 如果一个非空流是阻塞流，则空流可以阻塞该非空流中的操作
- 如果一个非空流是非阻塞流，则它不会阻塞空流中的操作
在下面的部分中，将介绍如何使用阻塞流和非阻塞流。
阻塞流和非阻塞流:
前头使用的cudsStreamCreate()创建的就是阻塞流
阻塞流和空流的相互依赖关系(同步关系)如下:
当操作被发布到空流中，在该操作被执行之前，CUDA上下文会等待所有先前的操作发布到所有的阻塞流中。此外，任何发布到阻塞流中的操作，会被挂起等待，直到空流中先前的操作执行结束才开始执行
这段有点绕, 但看下头的例子:
这段代码的结果是, 直到核函数kernel_1执行结束，kernel_2才会在GPU上开始执行，kernel_2执行结束后，kernel_3才开始执行
创建非阻塞流:
__host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t * pStream, unsigned int flags)
使用这个函数创建:
flags参数用于指定是否创建非阻塞流:
指定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
所以要求是调用这个函数前, 不能向指定的流添加任何任务
可配置事件:
这里具体还不是很清楚其他flags 的意义
CUDA运行时提供了一种方式来定制事件的行为和性能:
__host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t * event, unsigned int flags)
书里翻译的极其奇怪, 这里使用官方CUDA源文:
Creates an event object for the current device with the specified flags. Valid flags include:
- cudaEventDefault: Default event creation flag.
- cudaEventBlockingSync: Specifies that event should use blocking synchronization. A host thread that uses cudaEventSynchronize() to wait on an event created with this flag will block until the event actually completes.
- cudaEventDisableTiming: Specifies that the created event does not need to record timing data. Events created with this flag specified and the cudaEventBlockingSync flag not specified will provide the best performance when used with cudaStreamWaitEvent() and cudaEventQuery().
- cudaEventInterprocess: Specifies that the created event may be used as an interprocess event by cudaIpcGetEventHandle(). cudaEventInterprocess must be specified along with cudaEventDisableTiming.
-
default就是创建一个普通的事件
-
cudaEventBlockingSync创建一个会阻塞host 的event
即指定了所创建的event应该使用cudaEventSynchronize()等待事件完成, 而不是异步进行 -
cudaEventDisableTiming创建一个只用来同步, 而不用来计算时间的event
这样就能减少记录时间戳的消耗, 提高cuudaStreamWaitEvent和cudaEventQuery的调用性能
-
cudaEventInterprocess指定event可以被用来作为inter-process event
这玩意暂时用不到
6.2 并发内核执行:
这一节就是将上头讲到的几个例子实践一波:
包括以下几个方面：
-
使用深度优先或广度优先方法的调度工作
-
调整硬件工作队列
-
在Kepler设备和Fermi设备上避免虚假的依赖关系
-
检查默认流的阻塞行为
-
在非默认流之间添加依赖关系
-
检查资源使用是如何影响并发的
非空流中的并发内核
在本节中，将使用NVIDIA的可视化性能分析器（nvvp）可视化并发核函数执行
为了能更好的体现出差异性, 核函数在GPU上驻留的时间要足够长, 所以本例中使用了多个完全相同的核函数:
创建一组非空流, 并为每个流添加任务:
并加入到默认流中:
输出结果:
在Tesla K40上运行
Fermi GPU上的虚假依赖关系:
由于基本不可能使用到Fermi架构的GPU了, 所以这节可略
但是为了理解虚假依赖关系, 还是看看较好
这个时序图和上一节讲的相同, 仅有最后一个任务和另一个流的第一个任务能并行
避免虚假依赖关系:
使用OpenMP的调度操作:
本小节使用OpenMP创建多个主机线程进行调度
这里就是将上一个小节中的循环调度操作改用OpenMP进行多线程并行
可以看到这种简单的并行并没有带来明显的性能差异
仅有当主机&设备都有计算任务时, 主机上的多线程才能更加有效
用环境变量调整流行为
之前讲到Kepler设备支持的硬件工作队列的最大数量是32, 然而，默认情况下并发硬件连接的数量被限制为8
由于每个连接都需要额外的内存和资源，所以设置默认的限制为8，减少了不需要全部32个工作队列的应用程序的资源消耗
可以使用CUDA_DEVICE_MAX_CONNECTIONS环境变量来调整并行硬件连接的数量，对于Kepler设备而言，其上限是32
有几种设置该变量的方法:
-
Linux中使用Bash & Bourne Shell:
-
C-Shell:
-
直接在主机程序中进行设定:
修改上头的程序以呈现出流数量大于工作队列数量的情况:
图6-10展示了8个流，但是只有4路并发。因为现在只有4个设备连接，两个流共享一个队列
GPU资源的并发限制:
有限的内核资源可以抑制应用程序中可能出现的内核并发的数量
在之前的例子中，由于避免出现资源导致的并发限制, 所以启动内核时只有一个线程
在本例中, 将提高核函数的运行时配置, 以测试资源导致的并发限制
如图6-12所示，图中只实现了8路并发，即使CUDA设备连接的数量被设置为32。因为GPU无法分配足够的资源来执行所有符合条件的内核，所以并发性是有限的
所以可知, 即使工作队列的上限是32, 但是由于硬件设备资源的原因, 实际的最大并发流个数还是需要实际测试的
默认流的阻塞行为:
本小节中说明了默认流在非空流中是如何阻塞操作的
将上头例子中的主机程序修改如下:
因为第三个内核在默认流中被启动，所以在非空流上所有之后的操作都会被阻塞，直到默认流中的操作完成
创建流间依赖关系:
在理想情况下，流之间不应该有非计划之内的依赖关系（即虚假的依赖关系）
但在复杂的应用程序中，不同的流之间的工作很可能需要相互配合
本例就是利用上头的cudaStreamWaitEvent创建流之间的依赖关系
假设想让一个流中的工作在其他所有流中的工作都完成后才开始执行
而后将这些事件插入到其他流的末尾, 用以标志流的结束, 并使最后一个流等待这些事件的完成:
6.3 重叠内核执行 & 数据传输:
在前一节中，已经介绍了如何在多个流中并发执行多个内核。在本节中，将学习如何并发执行内核和数据传输
重叠内核和数据传输表现出不同的行为，并且需要考虑一些与并发内核执行相比不同的因素
Fermi GPU和Kepler GPU有两个复制引擎队列：
- 一个用于将数据传输到设备
- 另一个用于从设备中将数据提取出来
所以最多可以重叠两个不同方向的数据传输, 否则, 所有的数据传输都将是串行的
在应用程序中，还需要检验数据传输和内核执行之间的关系，从而可以区分以下两种情况：
-
如果一个内核使用数据A，那么对A进行数据传输必须要安排在内核启动前，且必须位于相同的流中
这种情况下，要实现数据传输和内核执行之间的重叠会更复杂，因为内核依赖数据作为输入
-
如果一个内核完全不使用数据A，那么内核执行和数据传输可以位于不同的流中
这种情况下，实现内核和数据传输的并发执行是很容易的：将它们放置在不同的流中，这就已经向运行时表示了并发地执行它们是安全的
使用深度优先调度重叠:
本例中将使用上头熟悉的向量加法核函数:
其中增大了n_repeat从而以延长内核的执行时间
实现向量加法的CUDA程序，其基本结构包含3个主要步骤：
-
将两个输入向量从主机复制到设备中
-
执行向量加法运算
-
将单一的输出向量从设备返回主机中
之前执行的操作是将全部的数据一次性拷贝到设备内存中, 而后在执行计算, 最后再将结果一次性拷贝回来
而现在将向量加法划分为多个子任务, 这样就能实现不同子任务的数据拷贝和数据计算的重叠
要重叠数据传输和内核执行，必须使用异步复制函数
接下来需要将n个子任务分配到n个流中:
现在，可以使用一个循环来为几个流同时调度iElem个元素的通信和计算
通过将数据传输和该数据上的计算放在同一个流中，输入向量、内核计算以及输出向量之间的依赖关系可以被保持
并且由于多个子任务互不影响, 所以可以分配到多个流中并行执行
为了进行对比，此例还使用了一个阻塞实现来计算基准性能：
图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进行的网格依赖性分析也可以帮助消除虚假的依赖关系
使用广度优先调度重叠:
修改核函数如下以实现广度优先的效果
图6-17显示了在K40设备上只使用一个硬件工作队列时的时间轴
与深度优先的方法相比它没有明显的差异，因为Kepler的双向调度机制有助于消除虚假的依赖关系
但如果在Fermi设备上运行相同的测试，在整体性能方面会发现，使用广度优先方法不如使用深度优先方法。由主机到设备复制队列上的争用导致的虚假依赖关系，在主机到设备间的传输完成前，将阻止所有的内核启动
所以结论是:
- 对于Kepler架构的GPU, 通常无需关注其工作调度顺序
- 对于Fermi设备, 则需要注意这些问题
6.4 重叠GPU和CPU的执行:
相对而言，实现GPU和CPU执行重叠是比较简单的，因为所有的内核启动在默认情况下都是异步的
只需简单地启动内核，并且立即在主机线程上实现有效操作，就能自动重叠GPU和CPU执行
本节的示例主要包括两个部分：
-
内核被调度到默认流中
-
等待GPU内核时执行主机计算
使用以下简单的内核实现一个向量与标量的加法：
本例中使用了3个CUDA操作（两个复制和一个内核启动）。记录一个停止事件，以标记所有CUDA操作的完成。
由于上头使用的全是异步操作, 所以在这之后控制权立即返回到主机, 此时即可进行计算操作:
这里做的工作就是计数器不断++
以下是在Tesla K40上使用nvprof的输出。在等待GPU操作完成时，主机线程执行14606次迭代。
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一次，只执行一次回调
- 回调函数执行时, 阻塞队列中排在其后面的工作，直到回调函数完成
小栗子:
总结:
流的概念是CUDA编程模型的一个基本组成部分。允许高级CUDA操作在独立的流中排队执行，CUDA流可以实现粗粒度并发。因为CUDA支持异步操作和大多数版本的运行时函数，所以它可以在多个CUDA流之间调度计算和通信。
从概念上讲，如果CUDA操作之间存在依赖关系，则它们必须在同一个流中被调度。例如，为了确保应用程序的准确无误，内核必须在同一个流中被调度，并在它使用的任何数据传输后进行。另外，没有依赖关系的操作可以在任意的流中被调度。在CUDA中，通常可以使用3种不同类型的重叠方案来隐藏计算或通信延迟：
-
在设备上重叠多个并发的内核
-
重叠带有传入或传出设备数据传输的CUDA内核
-
重叠CPU执行和GPU执行
为了充分利用设备，并确保最大的并发性，还需要注意以下问题：
-
平衡内核资源需求和并发资源需求。在设备上一次启动过多的计算任务，可能会导致内核串行，这会使得硬件资源的工作块变得可用。但是，也需要确保设备没有被充分利用，一直有工作在排队等待执行。
-
如果可能的话，避免使用默认流执行异步操作。放置在默认流中的操作可能会阻塞其他非默认CUDA流的进展。
-
在Fermi设备上，从深度优先和广度优先两方面考虑主机的调度。这个选择可以通过消除共享硬件工作队列上的虚假依赖关系，显著影响其性能。
-
要注意隐式同步的函数，并且充分利用它们和异步函数来避免性能的降低。
此外，本章还介绍了CUDA可视化性能分析器（nvvp）在可视化GPU执行中的作用。nvvp允许确认操作重叠的条件，并且易于多个流行为的可视化。
1万+
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
