cuda c权威编程指南笔记

一、基于CUDA的异构并行计算

1.1并行计算

并行性：包括任务并行和数据并行。

当多任务或函数可以独立的大规模的并行执行时，就是任务并行，任务并行的重点在于利用多核系统对数据进行分配。

当同时处理许多数据时，就是数据并行，数据并行的重点在于数据的分配。

CUDA编程非常适合解决数据并行计算的问题

数据划分方法：块划分和周期划分

块划分：每个线程计算一部分数据，通常这些数据有相同的大小

周期划分，每个线程计算数据的多部分

二、CUDA编程模型

关于cudaMalloc()参数的解释

int *da;
cudaMalloc((void**)&da,size);

这里da是int指针，在cudaMalloc()的参数中需要一个指向指针的指针（即void**），来将这个指针的值改变为GPU中的内存地址。如果直接传这个指针，只能改变这个指针指向地址的值而不能改变该指针的指向地址（值传递，实际上形参只复制了指针指向的地址）。

关于cudaMemcpy()的同步问题：

a）同一个stream（包括默认stream）中的kernel后面的copy函数总是会等到kernel结束才执行copy，无论是同步版的cudaMemcpy()还是异步版的cudaMemcpyAsync()。
b）对于host端而言，cudaMemcpy()是同步返回的，而cudaMemcpyAsync()是异步返回的（不等真正执行完就返回的）。但这并不影响和kernel之间的配合。

三、CUDA执行模型

通过了解CUDA的执行模型来优化指令吞吐量。

3.1CUDA执行模型概述

GPU是围绕着一个流式多处理器stream Multiprocessor（SM）的可拓展阵列搭建的，以Fermi架构为例，说明GPU的关键组件：

• CUDA核心
• 共享内存\一级缓存
• 寄存器文件
• 加载/存储单元
• 特殊功能单元
• 线程束调度器

下图为Ampere架构的A100GPU的SM架构：

block，grid是抽象概念，物理层次是SM，warp。当启动kernel时，block被分布在可用的SM上运行，block一旦被调度到一个SM上，其中的线程只会在那个指定的SM上并发执行。多个block可能被分配到同一个SM上。

CUDA采用SIMT单指令多线程架构来管理和执行线程，每32个为一组，称为warp。warp中所有线程同时执行相同的指令，每个线程都有自己的指令计数器和寄存器状态，因此每个线程会有独立的程序执行路径。（SIMT侧重于线程级并行，且多线程执行的是同一指令（同一kernel函数））

逻辑层次和物理层次对应关系如下图（左侧第二个应该为block）：

block和warp的对应关系如下图：

同一个block中的线程可以同步，但是block间线程无法同步。

由于SM的资源限制，活跃的warp是有限的，但是warp间的切换是没有开销的，因为硬件资源已经分配到了SM上的线程和块中。当warp因为任何理由闲置时（例如从设备内存中读取数值），SM可以从同一SM上的常驻线程块中调度其他可用warp。

3.2理解warp执行的本质

3.2.1 warp与block

从逻辑视角看，线程可以组织成一维二维三维的block，然而，从硬件的角度看，所有的线程被组织成一维的，然后将其分配到warp中。

block中warp的数量=向上取整【block中线程总数/warp大小（32）】

因此，当block中线程数不能被32整除时，最后的warp有些线程就不会活跃，但是它们仍会消耗SM的资源比如寄存器。

3.2.2 warp分化

warp分化是指同一warp的线程在执行有分支的程序时，warp将连续连续执行每一个分支路径，同时会禁用不执行这一路径的线程。这会导致并行线程的数量降低，导致性能明显下降，且条件分支越多，并行性削减越明显。

造成warp分化的原因是SIMT特性，同一warp中的线程必须同时执行同样的指令。且GPU没有复杂的分支预测机制，因此必须通过执行每一个分支的方式完成分支控制。

warp对性能影响很大，因此我们必须尽可能避免warp分化。所依据的原理就是根据block中线程的warp分配是确定的，我们可以通过确保同一warp中的线程执行的分支相同，来避免warp分化。

举个简单的例子：

上面这个例子就将控制分支的分支粒度由1变为wrap大小

3.2.3 资源分配

warp的本地执行上下文有以下资源组成

• 程序计数器
• 寄存器
• 共享内存

SM处理的每个warp的执行上下文在其声明周期内是保存在硬件上的，上下文切换没有损失，但这也导致了每个SM可处理的warp是有限的，warp线程消耗寄存器越多，SM可处理的warp就越少，一个block消耗的共享内存越多，SM中可以被同时处理的线程块就越少。

warp分为三种类型：

• 选定的warp：正在执行的warp是选定的
• 阻塞的warp：没有做好准备执行
• 符合条件的warp：准备执行但尚未执行

3.2.4 隐藏延迟

指令延迟：在指令发出和完成之间的时钟周期被定义为指令延迟

当每个时钟周期中所有的线程调度器都有一个符合条件的warp时，可以达到计算资源的完全利用。

延迟隐藏不是让计算单元的延迟消失，而是让warp调度器不受指令延迟影响，GPU的指令延迟被其他warp的计算隐藏。

利特尔原则：所需warp数=延迟*吞吐量

例如：内核中一条指令的平均延迟是5个周期。为了在每个周期内执行6个warp的吞吐量，则至少需要30个未完成的warp。

3.2.5 占有率

占有率：每个SM中活跃的warp数占最大warp数量的比值

占有率=活跃线程束数量/最大线程束数量

最大线程束数量由硬件决定，比如3090就是48

block和grid大小的准则：

• 每个block中线程数量是warp大小的倍数
• 避免block太小，至少为128或者256
• block的数量要远远大于SM的数量，从而保证足够的并行性

3.2.6 同步

系统级同步：cudaDeviceSynchronize()

block级同步：__device__ void syncthread(void)

这个函数会使同一block中的线程都达到这个同步点再继续执行。由于强制线程空闲，对性能有不利影响。

需要同步的情况：

block内的线程需要同步，当使用共享内存时，可能会出现race condition，这时利用同步需要规定好读写顺序。

block间线程同步没有办法，只能通过cudaDeviceSynchronize()进行block同步。

3.2.7可拓展性

简单的来说，可拓展性就是指当增加计算单元时性能也可以跟着提升。

CUDA内核启动时，block分布在多个SM中，block可以并行或连续或任意的顺序执行，这种独立性使CUDA程序在任意数量的计算核心间可以拓展

3.3 并行性的表现

减少block的维数增加block的个数会增加并行性

block的x维应该是warp大小的倍数

并行性、占有率、内存吞吐量等指标都不能单独决定性能

3.4 避免分支分化

避免分支分化的原理就是3.2.2里讲的尽量使同一个warp中的线程执行相同的控制分支。本节以规约问题为例介绍怎么避免分支分化。

reduction问题是经典的并行计算问题，基本思想如下：

问题描述：

若要使用并行算法来完成，可以通过相邻配对和交错配对的方式来完成。

在GPU中实现并行规约算法，由于线程块间不能同步，所有在block中并行规约结果后需要在host串行将结果相加起来。

并行核函数：

原理图：

main函数：

这是相邻配对的方式，每个warp中有一半的线程会分支分化，我们对此进行优化。

将其改为

相比于第一种规约，只有前面一半\四分之一\八分之一\。。。的线程进行计算，因此避免了warp分化。

再将其改为交错配对

三种核函数性能对比：

发现第三个比第一个快了1.69倍，比第二个快了1.34倍，与第二个相比，性能提升的原因是reduceInterleaved函数里的全局内存加载/存储模式导致的，在第4章里会介绍更多有关于全局内存加载/存储模式对内核性能的影响。

3.5展开循环

在CUDA中，循环展开的意义重大，通过减少指令消耗和增加更多的独立调度指令来提高性能。因此，更多的并发操作被添加到流水线上，以产生更高的指令和内存带宽。这为线程束调度器提供更多符合条件的线程束，它们可以帮助隐藏指令或内存延迟。

3.5.1 展开规约

在第四节的规约中，每个线程对应一个数据，每个block对应一个数据块。现在每个block展开两个数据块的处理（问题是原来不是并行的吗，这相当于循环吗？展开循环避免了许多指令和减少了计数，这里也会较少许多指令）

这条语句将每个block添加了相邻block的元素

速度快了3.42倍，再进一步展开，将其展开为4、8个数据块

3.5.2 展开线程的规约

在循环处理数据时，活跃线程数不断减半，当活跃线程减为32，就不会再较少了，这部分可以从循环中展开。

注意临时指针是volatile修饰的，它告诉百年一起每次赋值时需要将vmem的值存回全局内存，并且从全局内存读取，保证编译器不会对全局或共享内存优化读写。(为什么这个要，但是前面却不需要？？)

修改后的核函数如下

3.5.3 完全展开的规约

如果编译时已知一个循环中的迭代次数，就可以把循环完全展开。因为在Fermi或Kepler架构中，每个块的最大线程数都是1 024（参见表3-2），并且在这些归约核函数中循环迭代次数是基于一个线程块维度的，所以完全展开归约循环是可能的：

内核时间再次有了小小的改善，它的执行比reduceUnrollWarps8快1.06倍，比原来的实现快9.16倍

3.5.4 模板函数的规约

虽然可以手动展开循环，但是使用模板函数有助于进一步减少分支消耗。在设备函数上CUDA支持模板参数。如下所示，可以指定块的大小作为模板函数的参数：

相比reduceCompleteUnrollWarps8，唯一的区别是使用了模板参数替换了块大小。检查块大小的if语句将在编译时被评估，如果这一条件为false，那么编译时它将会被删除，使得内循环更有效率。

该核函数一定要在switch-case结构中被调用。这允许编译器为特定的线程块大小自动优化代码，但这也意味着它只对在特定块大小下启动reduceCompleteUnroll有效：

3.5.4 总结

注意，最大的相对性能增益是通过reduceUnrolling8核函数获得的，在这个函数之中每个线程在归约前处理8个数据块。有了8个独立的内存访问，可以更好地让内存带宽饱和及隐藏加载/存储延迟。可以使用以下命令检测内存加载/存储效率指标：

表3-6总结了所有核函数的结果。在第4章，将会更加详细地介绍全局内存访问，并且会对内存访问如何影响内核性能有更深的了解。

问题回到了隐藏延迟：为什么在warp中大量独立的内存操作可以隐藏延迟？？

3.6 动态并行

GPU动态并行允许在GPU端直接创建和同步新的GPU内核，有以下几点优点：

• 动态并行提供了一个更有层次结构的方法，在这个方法中，并发性可以在kernel中的多个级别表现出来。
• 有了动态并行，可以推迟到运行时决定需要在GPU上创建多少个block和grid，可以动态的利用GPU硬件调度器和加载平衡器，并进行调整以适应数据驱动或工作负载
• 在GPU端创建kernel减少了host和device之间的控制和数据传输。

本节以动态并行实现递归规约为例介绍动态并行。

3.6.1 嵌套执行

• 避免大量嵌套有利于提升性能。
• 同步对性能和正确性至关重要，但是减少block内部的同步次数会使嵌套内核的效率更高。
• 在每一个嵌套层上设备运行时都要保存额外的内存，所以内核嵌套的最大数量可能是受限制的。

四、全局内存

4.1 CUDA内存模型概述

对于程序员来说，有两种类型存储器：可编程与不可编程

不可编程：一级缓存和二级缓存

可编程：

• 寄存器
• 本地内存
• 共享内存
• 全局内存
• 常量内存
• 纹理内存

每个kernel中的每个线程都有自己的寄存器和本地内存，block内线程共享共享内存，device中的所有线程都可访问全局内存，所有线程都可访问只读的常量内存和纹理内存。

4.1.1 寄存器

寄存器式GPU上最快的存储器，kernel中没有其他修饰符修饰的变量通常存在寄存器中，kernel声明的数组中，

如果引用该数组的索引是常量且能够在编译时确定，那么该数组也存储在寄存器中。

寄存器变量对线程私有，声明周期与kernel的声明周期相同。

若kernel使用了超过硬件限制数量的寄存器，则会用本地内存代替多占用的寄存器。这种寄存器溢出会对性能带来不利影响。

我们可以显式的为kernel加上额外信息来帮助编译器优化。

限制了每个SM最大block和每个block最多线程

4.1.2 本地内存

除了寄存器溢出的变量会存储在寄存器上，编译器可能存放在本地内存中的变量有：

• 在编译时使用未知索引引用的本地数组
• 可能会占用大量寄存器空间的较大本地结构体或数组
• 其他不满足kernel寄存器限定条件的变量

需要注意的是，本地内存本质上与全局内存在同一块存储区域，因此高延迟，低带宽。

4.1.3 共享内存

当用__shared__修饰的变量存放在共享内存中。

相比本地内存和全局内存，具有更高的带宽和更低的延迟

共享内存被block内线程共享，因此要注意同步问题，使用__syncthreads进行同步。

SM中的一级缓存和共享内存都使用64KB的片上内存，是静态划分的，但是可以在运行时动态配置

4.1.4 常量内存

常量变量用__constant__修饰，必须在全局空间内和所有核函数外进行声明。其必须通过下面函数进行初始化。

大多数情况下这个函数是隐式同步的

4.1.5 纹理内存

4.1.6 全局内存

全局内存可以被静态或动态分配，静态通过__device__来分配内存，动态通过cudaMalloc()和cudaFree()来分配和释放全局内存。

访问全局内存也要注意同步问题，多个线程并发的修改内存的同一位置会导致未定义程序行为。

内存对齐：全局内存常驻于device内存中，可以通过32字节、64字节128字节的内存事务进行访问，这些内存事务必须自然对齐，也就是说首地址必须是32字节，64字节或128字节的倍数。

4.1.7 GPU缓存

GPU cache是不可编程的，在GPU上有四种缓存：

• 一级缓存
• 二级缓存
• 只读常量缓存
• 只读纹理缓存

每个SM有一个一级缓存，所有的SM共享一个二级缓存，一级缓存和二级缓存都能够用来存储本地内存和全局内存中的数据，包括寄存器溢出的部分。每个SM有一个只读常量缓存和只读纹理缓存。

L1cache和局部内存是同一块存储区域。L2cache被所有SM共享但是其速度要快于全局内存

在GPU中，只有内存加载操作可以被缓存，内存存储操作不能被缓存

4.1.8 CUDA变量声明总结

4.1.9 静态全局内存

4.1.6说到可以通过__device__声明静态全局内存，关于静态全局内存的使用有几点需要注意的

通过__device__声明的变量只是一个符号，不能通过这个变量访问GPU中的全局内存变量。即使在同一文件中可见，主机代码也不能访问设备变量，设备代码也不能访问主机变量。

要想访问GPU的全局内存变量，应该通过cudaMemcpyToSambol()和cudaMemcpyFromSambol()访问。

另外还可以cudaGetSymbolAddress()获得全局变量的统一虚拟地址（UAV），注意不能用&符号获取地址。

4.2 内存管理

尽可能减少主机与设备之间的传输

4.2.3 固定内存

因为虚拟内存的原因，在host端分配的内存是可分页的（pageable），可能在虚拟内存中，当GPU要访问主机端的数据时就可能发生page fault需要页面置换，从而降低效率。

我们可以使用cudaMallocHost()函数来分配host内存，这样这些内存的页面是锁定的即不会被置入swap区，从而提高了读写带宽。通过cudaFreeHost()来释放主机内存。

主机与设备间的内存传输

• 与分页内存相比，固定内存分配和释放成本更高，但是它为大规模数据传输提供了更高的传输吞吐量
• 将许多小的传输批处理为一个更大的传输能提高性能，因为它减少了单位传输消耗
• 主机与设备之间的数据传输有时可以与内核执行重叠。第六章会详细讲解
• 应当尽量减少或重叠主机与设备间的数据传输

4.2.4 零拷贝内存

之前讲过host不能直接访问device变量，device不能直接访问host变量，但是有一个例外：即零拷贝内存（zero copy mem）

零拷贝内存的优势：

• 当设备内存不足时可以使用主机内存
• 避免主机与设备之间的显示传输（使用零拷贝内存则会自动隐式传输）
• 提高PCLe的传输率

使用cudaHostAlloc()来给零拷贝内存分配空间

flag有四种：

可以使用cudaHostGetDevicePointer()函数获取设备端的指针

总结一下，零拷贝内存就是将显示的内存传输改为隐式，减少了编程者的工作量，对于少量数据来说零拷贝内存是一个不错的选择，但是对于由PCLe总线连接的离散GPU上的更大数据集来说，零拷贝内存不是一个好选择，它会导致性能显著下降。

有两种常见的异构计算架构：集成架构和离散架构（集成显卡和独立显卡）

在集成架构中，CPU和GPU集成在一个芯片上，并在物理地址上共享内存。在这种架构中，由于无需再PCLe总线上备份，所以零拷贝内存在性能上和可编程性方面可能更佳。

对于通过PCLe总线连接的离散系统而言，零拷贝内存只在特殊情况下有优势。

另外特写需要注意的，零拷贝内存被device和host共享，要注意同步问题。

4.2.5 统一虚拟寻址

统一虚拟寻址（UVA）指主机内存和设备内存共享同一个内存空间

UVA可以在零拷贝内存的基础上更进一步解放程序员，使用零拷贝内存时还需要创建host指针和device指针两个指针，有UVA后不需要两个指针，只需要一个指针就可以被host和device访问。

对比零拷贝内存的代码和UVA的代码

有了UVA可以直接将指针传给核函数

4.2.6 统一内存寻址

在CUDA 6.0引入统一内存寻址,进一步简化了内存管理。其原理为创建了一个托管内存池，内存池已分配的空间可以用相同的内存地址在CPU和GPU上进行访问。

统一内存寻址和UVA不同，UVA只是创建了统一的虚拟内存空间，但是不会自动将数据从一个物理位置转移到另一个物理位置。它应用都是在主机端分配内存，会受到PCLe传输的影响，核函数延迟高。而统一内存寻址将内存和执行空间分离，因此可以根据需要将数据透明的传输到主机或设备上，以提高局部性和性能。

可以通过__managed__来静态声明一个托管变量，但是只能在文件范围和全局范围内进行

或者通过cudaMallocManaged()来动态分配托管内存。

4.3 内存访问模式

核函数的内存请求是在DRAM(片外全局内存)和片上内存间以128字节或32字节的内存事务实现的。

一级缓存的缓存行是128字节。二级缓存的缓存行是32字节。可以在编译时选择是否开启一级缓存，默认开启,通过-Xptxas -dlcm-cg关闭缓存，-Xptxas -dlcm-ca开启。

4.3.1 对齐和合并访问

从全局内存读取数据有两个特性：对齐访问，全局访问

对齐访问指从内存读取到缓存时，必须从缓存行大小的整数倍地址开始读，比如一级缓存，每次读取的地址必须是128的倍数。

合并访问：CUDA模型的显著特征之一是指令必须以warp为单位进行发表和执行，存储也是一样。当warp的32个线程访问一个连续的内存块时，就会出现合并内存访问。

如何优化内存事务效率：用最少的事务请求满足最多的内存请求，即尽可能的减少访存次数。

4.3.2 全局内存读取

启用一级缓存的内存加载以128字节粒度进行加载，不启用则以32字节进行加载。在未对其的情况下，不启用缓存会使加载效率得到提升，因为一次加载的字节数较少，无用的数据部分会减少。但是缓存可以减少重复加载。

CPU与GPU一级缓存的区别：

CPU一级缓存优化了时间和空间局部性，GPU专为空间局部性设计，频繁访问一个一级缓存的内存位置不会增加数据留在缓存中的概率。

4.3.3 全局内存存储

内存写入只通过二级缓存，在32字节的粒度上被执行。内存事务可以同时被分为一段、两段和四段、

对齐存储效率显著高于非对齐存储。

4.3.4 结构体数组和数组结构体

结构体数组（SoA）数组结构体（AoS）是两种常见的数据组织方式，

struct innerStruct{
  float x;
  float y;
}

struct innerStruct myAos[n]

struct innerArray{
	float x[n];
  float y[n];
}

两种方式在内存中的结构如图所示：

使用SOA模式可以充分利用GPU的内存带宽，由于相同字段元素相邻存储，不仅可以合并内存访问，还可以对全局内存实现更高效的利用。

许多并行编程范式，尤其是SIMD类型范式更倾向于SOA，CUDA也倾向于使用SOA。

4.3.5 性能调整

展开技术

在第三章已经讲过，增加每个线程执行独立内存操作的数量可以提高性能。

对于IO密集型的核函数，内存访问并行有很高的优先级。

增大并行性

增大并行性主要通过修改核函数的配置实现，通过减少block内线程数，增加block的数量来增大并行性。但是block中线程的数量也不能太少，原因如下：

最大化带宽利用率

总结下要最大化贷款的利用率要从两个角度出发
一是提高DRAM和SM片上内存的有效内存的移动，避免浪费。并且要保证内存访问应当是对齐和合并的
二是提高并发内存操作，主要途径有（1）展开（2）修改核函数启动配置来提高并行性

4.4 核函数可达到的带宽

4.4.1 理论带宽与有效带宽

理论带宽：当前硬件可实现的绝对最大带宽
有效带宽：（读字节数+写字节数）* 10^-9/运行时间

4.4.2 矩阵转置问题

简单的矩阵转置代码：

void transposeHost(float *out, float *in, const int nx, const int ny) {
	for (int iy = 0; iy < ny; ++iy) {
		for (int ix = 0; ix < nx; ++ix) {
			out[ix*ny+iy] = in[iy*nx+ix];
		}
	}
}

矩阵转置有两种基本方法：行读取列存储，列读取行存储，其中行操作都是合并存储的，列操作都是交叉存储的。

如果禁用一级缓存，这两种方法结果相同，但是如果启用一级缓存，列读取行存储的有效带宽更高，因为读取时虽然是交叉读取，但是因为有缓存可以减少访存；但是行读取列存储时由于存储不经过一级缓存，所以缓存对其没有意义。

展开转置

将转置操作展开能获得更高的带宽

__global__ void transposeUnroll4Col(float *out, float *in, const int nx,
const int ny) {
  unsigned int ix = blockDim.x * blockIdx.x*4 + threadIdx.x;
  unsigned int iy = blockDim.y * blockIdx.y + threadIdx.y;
  unsigned int ti = iy*nx + ix; // access in rows
  unsigned int to = ix*ny + iy; // access in columns
  if (ix+3*blockDim.x < nx && iy < ny) {
    out[ti] = in[to];
    out[ti + blockDim.x] = in[to+ blockDim.x*ny];
    out[ti + 2*blockDim.x] = in[to+ 2*blockDim.x*ny];
    out[ti + 3*blockDim.x] = in[to+ 3*blockDim.x*ny];
  }
}

对角转置

使用瘦块（thin block）增加并行性

五、共享和常量内存

5.1 CUDA共享内存概述

5.1.1 共享内存

共享内存（SMEM）与一级缓存在片上，相比于全局内存，其延迟低20-30倍，带宽高大约10倍。

共享内存被block内的线程共享，当warp读取共享内存时，理想情况下每个请求应该在一个事务中完成，在最坏情况下，一个请求在32个事务中顺序执行。因此要避免多次共享内存请求。当多个线程访问共享内存中的同一个字，一个线程读取该字后，通过多播把他发送给其他线程。

共享内存被SM所以常驻线程共享，所以共享内存是限制设备并行性的关键资源。一个核函数使用的共享内存越多，处于并发活跃状态的线程块就越少。

可编程管理的缓存

在编写CPU程序时，缓存对于程序是透明的，我们不能直接操纵缓存，只能通过循环转换等方法优化缓存。循环转换是一种常用的缓存优化方法，通过重新安排迭代顺序，提高缓存的局部性。

而共享内存是可编程管理的缓存，我们可以通过在数据布局上提供更多细粒度控制和改善片上数据的移动，使得应用程序代码优化变得简单。

5.1.2 共享内存分配

共享内存可以动态或静态分配，其作用域可以分配为全局或局部。

通过__shared__修饰符声明变量为静态分配共享变量，如果在核函数内进行声明则是局部变量，如果在核函数外进行声明即为全局变量。

例如

__shared__ float tile[size_y][size_x];

若共享内存大小在编译时未知，可以动态声明，使用extern关键字声明未知大小的数组，在核函数调用时，将所需的字节数作为三重括号内的第三个参数

extern __shared__ int tile[];
kernel<<<grid,block,isize* sizeof(int)>>>

注意：只能动态声明一维数组！

5.1.3 共享内存存储体(bank)和访问模式

存储体（bank）

为了获得高内存带宽，共享内存被分为32个同样大小的内存模型,它们被称为存储体，它们可以被同时访问。有32个存储体是因为在一个线程束中有32个线程。共享内存是一个一维地址空间。根据GPU的计算能力，共享内存的地址在不同模式下会映射到不同的存储体中(稍后详述)。如果通过线程束发布共享内存加载或存储操作，且在每个存储体上只访问不多于一个的内存地址，那么该操作可由一个内存事务来完成。否则，该操作由多个内存事务来完成，这样就降低了内存带宽的利用率。

存储体冲突(bank conflict)

在共享内存中当多个地址请求落在相同的内存存储体中时，就会发生存储体冲突，这会导致请求被重复执行。

当线程束发出共享内存请求时,有以下3种典型的模式:

• 并行访问:多个地址访问多个存储体

• 串行访问:多个地址访问同一个存储体

• 广播访问:单一地址读取单一存储体

并行访问是最常见的模式，它是被一个线程束访问的多个地址落在多个存储体中。这种模式意味着，如果不是所有的地址，那么至少有一些地址可以在一个单一的内存事务中被服务。最佳情况是，当每个地址都位于一个单独的存储体中时，执行无冲突的共享内存访问。串行访问是最坏的模式，当多个地址属于同一个存储体时，必须以串行的方式进行请求。如果线程束中32个线程全都访问同一存储体中不同的内存地址，那么将需要32个内存事务，并且满足这些访问所消耗的时间是单一请求的32倍。
在广播访问的情况下，线程束中所有的线程都读取同一存储体中相同的地址。若一个内存事务被执行，那么被访问的字就会被广播到所有请求的线程中。虽然一个单一的内存事务只需要一个广播访问，但是因为只有一小部分字节被读取，所以带宽利用率很差。

若每个线程访问一个存储体，这是最优的并行访问模式，若多个线程访问一个存储体，有两种可能的情况：

• 如果线程访问同一个存储体中相同的地址，广播访问无冲突
• 如果线程访问同一个存储体中不同的地址，会发生存储体冲突

访问模式

内存存储体的宽度（字长）随计算能力的不同而变化：计算能力2.x的为四字节，计算能力3.x的为8字节

字长为四字节的存储模式：

字长为8字节的存储模式

内存填充

内存填充是避免存储体冲突的一种方法，以下图为例，若bank0发生大量存储体冲突,可以通过填充字的方式来避免

填充的内存不能用于数据存储，这样block中可用的共享内存会减少

5.1.4 配置共享内存

SM上的片上内存被共享内存和一级缓存共享。有两种配置方法：

• 按设备进行配置
• 按核函数进行配置

具体配置方法略去，不同计算能力的设备的片上内存不同，配置方法不同。

5.1.5 同步

共享内存的同步有两个基本方法

• 障碍（barriers）
• 内存栅栏（memory fences）

弱排序内存模型

CUDA采用弱排序内存模型，这意味着，内存访问不一定按照它们在程序中的顺序执行，为了显式的确定内存访问的顺序，必须通过barriers和memory fences来保证。

显示障碍（barriers）

void __syncthreads()作为显示障碍保证当所有线程都到达该点时才能继续执行

需要注意的是在条件语句中调用void __syncthreads()，很可能块中的线程无法到达相同的障碍点

if (threadID % 2 == 0) {
__syncthreads();
} else {
__syncthreads();
}

另外，前面已经说过block会以并行、串行的方式执行，这保证了CUDA是可拓展的，但这也会导致block间无法同步，若想实现block间的同步，可以在核函数中启动多个核函数并使用__syncthreads()进行同步。

内存栅栏（memory fences）

什么是内存栅栏（内存屏障，memory barriers）

在并行系统中，由于指令的乱序执行，实际的访存可能会和程序中的访存顺序不同，这就会导致程序不按照预想的形式执行，为了避免这种问题，引入内存栅栏，保证访存是按照程序中的顺序。

有三种范围的内存栅栏：block,grid,system

void __threadfence_block();

保证栅栏前该线程对所有共享内存和全局内存的写操作对同一块中的其他线程是可见的，即保证其他线程都知道该线程执行了这些访存操作

void __threadfence();

范围变为grid

void __threadfence_system();

范围变为system，全部设备

volatile

volatile的作用与c语言中相同，可以防止编译器优化，保证volatile修饰的变量都将存到全局内存中，避免缓存在寄存器和本地内存中。

5.2 共享内存的数据布局

介绍了方形和矩形两种常见的数据布局，以及其对应的矩阵转置算法，还有通过内存填充来避免存储体冲突。以行主序写列主序读在读的时候会出现存储体冲突，通过内存填充可以避免存储体冲突。

要区分在全局内存和共享内存矩阵转置的区别，全局内存要避免的是交叉访问尽量合并访问，共享内存要避免的是存储体冲突。

具体见原书代码吧，应当学会矩阵的转置以及内存填充。

5.3 减少全局内存的使用

本节以第三章提到的规约函数为例，介绍如何使用共享内存作为可编程管理缓存以减少全局内存的使用。

通过将规约函数中的数据处理放到共享内存中，减少了全局内存的使用，提高了性能。具体应当查看原书代码，学会使用局部内存。

5.4 合并的全局内存访问

5.5 常量内存

常量内存位于DRAM上，且有专门的片上缓存。常量内存主要有两种用途：

• 只读数据
• 当线程束中线程访问相同位置时

常量内存在设备端是只读的，在host是可读可写的

常量内存的访问模式不同于其他内存，当线程束中所有线程都访问相同的位置是，这个访问模式是最优的，当线程束访问不同的地址时，该访问就需要串行。

声明：通过__constant__声明

初始化：通过cudaMemcpySymbol(const void *symbol,const void *src,size_t count,size_t offset,cudaMemcpyKind kind),kind可省略，省略就表示默认cudaMemcpyHostDevice

与只读内存的比较

常量缓存与只读缓存

在设备上只读缓存和常量缓存都是只读的

每个SM资源有限，只读缓存48KB，常量缓存64KB

常量缓存在统一读取中可以更好运行，只读缓存更适合分散读取

声明与初始化（主机端）：

const float h_coef[] = {a0, a1, a2, a3, a4};
cudaMalloc((float**)&d_coef, (RADIUS + 1) * sizeof(float));
cudaMemcpy(d_coef, h_coef, (RADIUS + 1) * sizeof(float), cudaMemcpyHostToDevice);

设备端

__log() 例如：output[idx] += __ldg(&input[idx])强制使用只读缓存

5.6 warp洗牌指令

洗牌指令：只要两个线程在相同的线程束中，就允许这两个线程直接读取另一个线程的寄存器

首先必须要介绍束内线程（lane）的概念，每个warp中的束内线程有其独一无二的束内线程索引（laneID），每个warp有其线程束索引（warpID）
$$
laneID=threadID.x%32\
warpID=threadID.x/32
$$

对于二维线程块可以将其转为一维线程索引再求laneID和warpID

六、流与并发

一般来说，在CUDA中有两种级别的并发：

• 内核并发
• 网格并发

本章主要介绍网格并发以及如何用NVVP将内核并发执行可视化

6.1 流与事件

CUDA流目的：实现网格级并发（如何实现并发：通过异步操作，异步了如何确定执行顺序：通过流来确定执行顺序）

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

这些操作包括在主机与设备间进行数据传输，内核启动以及大多数由主机发起但由设备处理的其他命令。流中操作的执行相对于主机总是异步的。CUDA运行时决定何时可以在设备上执行操作。我们的任务是使用CUDA的API来确保一个异步操作在运行结果被使用之前可以完成。

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

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

6.1.1 CUDA流

流的两种类型：

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

如果没有显示指明流，那么数据传输和内核启动默认使用空流，前面几章的例子使用的都是空流。

接下来讲解如何创建销毁检查流：

声明：cudaStream_t stream;

创建：cudaStreamCreate(&stream);

销毁：cudaError_t cudaStreamDestroy(cudaStream_t stream);

当cudaStreamDestroy函数调用时，若流中有未完成的工作，该函数将立即返回，若所有工作已经完成，与流相关的资源将被自动释放。

因为CUDA流操作是异步的，不知道什么时候结束，所以有相关函数来检查流中操作是否以及完成

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

cudaStreamSynchronize强制阻塞主机，直到在给定流中所有的操作都完成了。cudaStreamQuery会检查流中所有操作是否都已经完成，但在它们完成前不会阻塞主机。当所有操作都完成时cudaStreamQuery函数会返回cudaSuccess，当一个或多个操作仍在执行或等待执行时返回cudaErrorNotReady。

在之前的代码中，常见的数据传输和执行内核操作比如：

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

在这个代码中，三个操作被发布到默认的流中，cudaMemcpy()主机同步等待数据传输完毕，在传输完成前，主机将强制空闲，内核启动是异步的，因此可以做到在内核启动后主机与设备的并行计算。

如果一直使用cudaMemcpy() 进行同步数据传输的话，主机端有无效的等待时间，也无法实现网格并行，可以通过cudaMemcpyAsync()来异步的数据传输，主机调用函数以后可以继续向前执行，由设备端继续执行，当然这样做必须三个函数的执行顺序必须按照发布顺序执行，因此需要通过显式的设置CUDA流来进行装载。

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

将流作为第五个参数。

需要注意的是，执行异步数据传输时，必须使用固定主机分页

cudaError_t cudaMallocHost(void **ptr, size_t size);
cudaError_t cudaHostAlloc(void **pHost, size_t size, unsigned int flags);

在非默认流启动核函数，必须将流作为内核执行配置的第四个参数

kernel_name<<<grid, block, sharedMemSize, stream>>>(argument list)

接下来在实例中看下非空流是如何实现网格并行的

for (int i = 0; i < nStreams; i++) {
  int offset = i * bytesPerStream;
  cudaMemcpyAsync(&d_a[offset], &a[offset], bytePerStream, 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]);
}

相比于之前的默认流代码，使用流的异步代码不必阻塞，在一个数据传输及核函数返回前可以启动下一数据传输和核函数，实现网格级并行，效果如下图。

6.1.2 流调度

虚假依赖关系

虽然fermi架构支持16路并发，但所有的流最终多路复用到单一的迎检工作队列，这就导致了虚假的依赖关系，阻碍了网格级并行

Hyper-Q

使用多个硬件工作队列，减少虚假依赖关系

6.1.3 流的优先级

计算能力3.5以上的设备，可以给流分配优先级

//分配优先级
cudaError_t cudaStreamCreateWithPriority(cudaStream_t* pStream, unsigned int flags,int priority);
//查询优先级
cudaError_t cudaDeviceGetStreamPriorityRange(int *leastPriority,int *greatestPriority);

cudaStreamCreateWithPriority这个函数创建了一个具有指定整数优先级的流，并在pStream中返回一个句柄。这个优先级是与pStream中的工作调度相关的。高优先级流的网格队列可以优先占有低优先级流已经执行的工作。流优先级不会影响数据传输操作，只对计算内核有影响。如果优先级超出了设备定义的范围，它会被自动限制为定义范围内的最低值或最高值。

cudaDeviceGetStreamPriorityRange的返回值放在leastPriority和greatestPriority中，如果当前设备不支持优先级，将返回0；

6.1.4 事件

CUDA事件与流中特定点相关联，可以执行两个基本任务：

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

声明：cudaEvent_t event;

初始化：cudaError_t cudaEventCreate(cudaEvent_t* event);

销毁：cudaError_t cudaEventDestroy(cudaEvent_t event);

在流中插入事件：cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);

等待一个事件结束:cudaError_t cudaEventSynchronize(cudaEvent_t event); 类似cudaStreamSynchronize

查询一个事件是否执行完：cudaError_t cudaEventQuery(cudaEvent_t event); 类似cudaStreamQuery

记录两个事件之间的时间：cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t stop); 两个事件不一定属于同一个流

6.1.5 流同步

前面已经讲过，为了实现网格并行，非默认流中的操作是非阻塞的，但是会有主机和设备需要同步的时候，也就是流同步。

阻塞流与非阻塞流

前面已经介绍过空流与非空流，空流也就是默认流是同步流，其操作（内存操作）会阻塞主机，而非空流是异步流，其操作不阻塞主机，非空流可以进一步分成以下两种流：

• 阻塞流：空流可以阻塞其操作
• 非阻塞流：不会阻塞空流中操作

没看懂，回头再回来看吧！！！

可配置事件

没看懂sad

隐式同步

CUDA中有两种设备-主机同步：隐式同步与显示同步，隐式同步比如cudaMemcpy，了解隐式同步很有意义，因为没有考虑隐式同步的话会导致想不到的性能下降。

隐式同步包括：

• 固定页主机内存分配
• 设备内存分配
• 设备内存初始化
• 同一个设备上两个地址之间的内存复制
• 一级缓存/共享内存配置的修改

显示同步

显示同步的几种方法：

-

• 同步设备：cudaDeviceSynchronize()，使主机等待设备相关的计算与通信完成
• 同步流：cudaStreamSynchronize()，使主机等待所有该流中的操作完成
• 同步事件：cudaEventSynchronize()，使主机等待事件完成
• 事件同步（可以跨流）：cudaStreamWaitEvent(cudaStream stream,cudaEvent_t event)，使一个流等待一个事件完成，这个事件可能不属于这个流，这样就可以实现跨流同步。

6.2 并发内核执行

6.2.1 非空流的并发执行

对于支持Hyper-Q的设备，使用多个非空流可以实现内核并发执行

6.2.2 fermiGPU的虚假依赖关系

然而在fermi架构中，由于不支持hyper-Q,所有流被多路复用的一个硬件工作队列，产生虚假依赖关系，无法并行

可以采用广度优先顺序，确保工作队列中任务来自不同的流

6.2.3 使用OpenMP的调度操作

之前的例子中都是使用一个线程启动多个内核，为了进一步提升性能可以使用多个主机线程将操作调度到多个流去。

omp_set_num_threads(n_streams);
#pragma omp parallel
{
  int i = omp_get_thread_num();
  kernel_1<<<grid, block, 0, streams[i]>>>();
  kernel_2<<<grid, block, 0, streams[i]
  kernel_3<<<grid, block, 0, streams[i]>>>();
	kernel_4<<<grid, block, 0, streams[i]>>>();
}>>>();

这个例子中使用openmp效果一般，当每个流在内核执行之前、期间或之后有额外的工作待完成，那么使用多线程调度流可以显著提高性能。