第2章 CUDA中的线程组织
2.1 C++语言开发过程
1. 用文本编辑器写一个源代码(source code)
2. 用编辑器对源代码进行预处理、编译、汇编并链接必要的文件得到可执行文件。这些步骤往往可由一个命令完成。
3. 运行可执行文件得到结果。
2.2 CUDA中的Hello World程序
2.2.1 只有主机的函数的CUDA程序
CUDA程序的编译器驱动nvcc在编译一个CUDA程序时，会将纯粹的C++代码交给C++编译器(如前面提到的g++或cl)去处理，它自己则负责编译剩下的部分。CUDA程序的源文件的扩展名是.cu。
2.2.2 使用核函数的CUDA程序
cuda中的核函数必须被限定词__global__修饰。核函数返回类型必须是空类型。
一个典型、简单的CUDA程序的结构具有以下形式
头文件包含
常量定义(或者宏定义)
C++ 自定义函数和CUDA核函数声明(原型)
C++ 自定义函数和CUDA核函数的定义(实现)
int main(void)
{
分配主机与设备内存;
初始化主机中的数据;
将某些数据从主机复制到设备;
调用核函数在设备中进行计算;
将某些数据从设备复制到主机;
释放主机与设备内存;
return 0;
}主机在调用一个核函数时，必须指明需要在设备中指派多少个线程。
2.3 CUDA中的线程组织
2.3.2 使用线程索引
核函数调用格式：
kernal<<<grid_size, block_size>>>();其中三尖括号中的数来指明核函数中的线程数目及排列情况。grid_size表示线程块的个数， block_size表示每个线程块中的线程数。
一个核函数若干线程(thread)构成一个线程块(block)。每个线程块含有相同数目的线程，该数目即为线程块大小(block size)。
一个核函数的全部线程块构成一个网格(grid)，而线程块的个数即为网格大小(grid size)。核函数中线程数目 = grid size * block size；
执行配置(execution configuration) : <<<grid_size, block_size>>>。一般来说，grid_szie与block_size是一个结构体类型(dim3 grid_size(x, y, z), dim3 block_size(x, y, z))的变量，但也可以是一个普通的整形变量， 没有指定的成员值将默认取1。
一个核函数虽然可以指派巨大数目的线程数目，但在执行时能够同时活跃的线程数目(不活跃的线程处于等待状态)是由硬件(主要是CUDA核心数，寄存器数目)和软件(即核函数中代码)决定的。
2.3.3 推广至多维网格
gridDim.x/y/z :该变量数值等于执行配置中变量grid__size.x/y/z的数值
blockDim.x/y/z :该变量数值等于执行配置中变量block__size.x/y/z的数值
blockIdx.x/y/z :该变量指定一个线程块在一个网格中x/y/z维度的线程块索引指标， 取值范围是0 ~~~~ gridDim.x/y/z - 1；
threadIdx.x/y/z :该变量指明一个线程在一个线程块中x/y/z维度的线程索引指标， 取值范围是0 ~~~~ blockDim.x/y/z - 1；
2.3.4 网格和线程块的大小限制
网格与线程块大小的限制
网格大小在x, y, z 三个方向最大允许值分别是2^31 - 1， 65535(2^16 - 1)， 65535;
线程块大小在x, y, z 三个方向最大允许值分别是1024，1064，64;
一个线程块中最大线程数blockDim.x * blockDim.y *blockDim.z<= 1024;
2.5 用nvcc编译CUDA程序
-arch = compute_XY 指定一个虚拟架构的计算能力；
-code = sm_ZW 指定一个真实架构的计算能力；
核函数中数据与线程对应: 用“单指令--多线程”的方式编写代码，顾客去掉循环，只需要将数组元素指标与线程指标一一对应即可。
unsigned int idx_x = blockDim.x * blockIdx.x + threadIdx.x;第3章 简单CUDA程序的基本框架
3.2.5 核函数要求
1. 核函数返回类型必须是void， 可以用return关键字，但不可返回任何值。
2. 必须使用限定符__global__。 也可以加上一些其他C++的限定符，如static，限定符次序可任意。
3. 函数名无特殊要求，而且支持C++中重载(overload)。
4. 不可支持可变数量的参数数量，即参数的个数必须确定。
5. 可以向核函数传递非指针变量，其内容对每个线程可见。
6. 除非使用统一内存编程机制，否则传给核函数的数组(指针)必须指向设备内存。
7. 核函数不可成为一个类的成员。通常做法是用一个包装函数调用核函数，而将包装函数定义为类的成员。
8. 在计算能力 3.5之前，核函数之间不能相互调用。计算能力3.5之后，引入动态并行(dynamic parallelism)机制，在核函数内部可以调用其他核函数。
9. if语言判定元素索引小于等于线程数目， 避免非法内存操作。
3.3.1 函数执行空间标识符
__global__ 修饰的函数称为核函数，主机调用，设备执行。若动态并行，则也可以在核函数调用。
__device__ 修饰的函数称为设备函数，只能被核函数或其他设备函数调用，在设备执行。
__host__ 修饰的函数就是主机端的普通C++函数。主机调用，主机执行。对于主机端函数，可省略该修饰符。
不能同时用__device__和__global__修饰同一个函数。
也不能同时用__host__和__global__修饰同一个函数。
可以用修饰符__forceinline__建议一个设备函数为内联函数。
第5章 获得GPU加速的关键
cuda程序获得高性能的必要条件
1. 数据传输比例较小；
2. 核函数的算数强度较高；
3. 核函数中定义的线程数目较多；
内建函数与数学函数使用，不需要包含任何额外的头文件。内建函数计算精度低，计算速度快。
第6章 CUDA的内存组织
6.1 CUDA的内存组织简介
现代计算机中的内存往往存在一种组织结构。该结构含有多种类型的内存，每种内存分别具有不同的容量和延迟。一般来说， 延迟低(速度高)的内存容量小，延迟高(速度低)的内存容量大。当前被处理的数据一般放于低延迟、低容量的内存中；当前没有被处理但之后将要处理的大量数据一般存放于高延迟、高容量的内存中。
6.2 CUDA中的不同内存
内存类型 | 物理位置 | 访问权限 | 可见范围 | 生命周期 | 延迟 | 带宽 |
全局内存 | 板上 | 可读可写 | 所有线程和主机端 | 由主机分配与释放 | 400~600个周期 | 200MB/s |
常量内存 | 板上 | 仅可读 | 所有线程和主机端 | 由主机分配与释放 | 400~600个周期 | 200MB/s |
纹理内存 | 板上 | 一般仅可读 | 所有线程和主机端 | 由主机分配与释放 | 400~600个周期 | 200MB/s |
寄存器内存 | 片上 | 可读可写 | 单个线程 | 所在线程 | 1个周期 | 8TB/s |
局部内存 | 板上 | 可读可写 | 单个线程 | 所在线程 | 400·600个周期 | 200MB/s |
共享内存/L1 缓存 | 片上 | 可读可写 | 单个线程块 | 所在线程块 | 1~32个周期 | 1.5TB/s |
6.2.1 全局内存
全局内存(global memory) : 核函数中所有的线程都能够访问其中的数据。 其容量基本就是显存容量。 其角色主要为核函数提供数据，并在主机与设备及设备与设备之间传递数据。
在处理逻辑上的二维或三维问题时，可以用cudaMallocPitch()和cudaMalloc3D()函数分配内存，用cudaMemcpy2D()和cudaMemcpy3D()复制数据，释放时候依然用cudaFree()函数；
6.2.2 常量内存
常量内存的访问速度比全局内存高，但前提是一个线程束中的线程要读取相同的常量内存数据；
除给核函数传递单个的变量外，还可以传递结构体；
6.2.3 纹理内存与表面内存
纹理内存(texture memory) 和表面内存(surface memory) 类似于常量内存，也是一种具有缓存的全局内存，有相同的可见范围和生命周期，而且一般仅可读（表面内存也可写）。不同的是，纹理内存和表面内存容量更大，而且使用方式和常量内存也不一样。
对于计算能力不小于3.5的GPU来说，对某些只读全局内存用__ldg()函数通过只读数据缓存(read-only data cache)读取，即可达到使用纹理内存的加速效果。对于Pascal及之后的架构来说，全局内存的读取在默认情况下使用了__ldg()函数。
6.2.4 寄存器
在核函数中不加任何限定符的变量一般来说就存放在寄存器中(register)，前提是寄存器内存足够能容纳该变量。
内建变量(gridDim, blockDim, blockIdx, threadIdx)都保存在特殊的寄存器内。
寄存器变量仅仅被一个线程可见，每一个线程都有一个变量n的副本。虽然在核函数的代码用了同一个变量名，但是不同线程中该寄存器变量的值是可以不同的。每个线程都只能对它的副本进行读写。
一个寄存器占有4字节的内存。一个双精度浮点数将使用两个寄存器。
6.2.5 局部内存
寄存器放不下的变量，以及索引值不能再编译时就确定的数组，都有可能放在局部内存中。
虽然局部内存用法上类似寄存器，但从硬件来看，局部内存只是全局内存的一部分。所以，局部内存延迟也很高。每个线程最多只能使用高达512KB的局部内存。
不同于寄存器，共享内存对整个线程块是可见。每个线程块拥有一个共享内存变量的副本。共享内存变量的值在不同的线程块可以不同。一个线程块中的所有线程都可以访问该线程块的共享内存变量副本，但是不能访问其他线程块的共享内存变量副本。
6.2.6 共享内存
共享内存的主要作用是减少对全局内存访问，或者改善对全局内存的访问模式。
计算能力 | 3.5 | 6.0 | 7.0 | 7.5 |
GPU代表 | K40 | P100 | V100 | RTX 2080 |
SM寄存器上限 | 64K | 64K | 64K | 64K |
单个线程块寄存器上限 | 64K | 64K | 64K | 64K |
单个线程寄存器上限 | 255 | 255 | 255 | 255 |
SM共享内存上限/KB | 48 | 64 | 96 | 64 |
单个线程块共享内存上/KB | 48 | 48 | 96 | 64 |
6.2.7 L1和L2缓存
从硬件结构角度，开普勒架构中的L1缓存和共享内存使用同一块物理芯片；在麦克斯韦架构和帕斯卡架构中，L1缓存和纹理缓存统一起来，而共享内存是独立的；在伏特和图灵架构中，L1缓存、纹理内存、共享内存三者统一管理。
6.3 SM及其占有率
6.3.1 SM的构成
一个GPU是由多个SM组成的。一个SM包含：
一定数量的寄存器；
一定数量的共享内存；
常量内存的缓存；
纹理内存和表面内存的缓存；
L1缓存；
2个（计算能力4.0）或4个（计算能力6.0）线程束调度器，用于在不同线程的上下文之间迅速地切换，以及为准备就绪的线程束发出执行指令。
执行核心：
若干整数型运算的核心
若干单精度浮点数运算的核心；
若干双精度浮点数运算的核心；
若干单精度浮点数超越函数的特殊函数单元；
若干混合精度的张量核心，有伏特架构引进，适用于机器学习；
6.3.2 SM的占有率
SM中驻留的线程数目除以SM可调度的最大线程数数目为SM占有率。
一个SM中最多能拥有的线程块个数Nb = 16(开普勒架构和图灵架构)或者Nb = 32(麦克斯韦架构，帕斯卡架构和伏特架构)。
一个SM中最多能拥有的线程个数Nt = 2048（c从开普勒架构到伏特架构） 或者N他= 1024（图灵架构）。
SM中线程的执行是以线程束为单位的，所以最好将线程块大小取为线程束32的整数倍。
任何不小于Nt/Nb 且能整除Nt的线程块大小都能得到100%的占有率。
有限寄存器数量对占有率的约束情况。若一个SM最多使用的寄存器数目为64k，如果希望常驻SM的线程2048个，核函数中每个线程最多使用32个寄存器。
有限共享内存对占有率的约束情况。
值得一提的是，用编译器选项--ptxas-options=-v可以报道每个核函数的寄存器使用数量。
cuda还提供了核函数的__launch_bounds__()修饰符核--maxregcount = 编译选项来让用户分别对一个核函数核所有核函数中寄存器的使用数量精选控制。
第7章 全局内存的合理使用
7.1 全局内存的合并有非合并访问。
在启用L1缓存的情况下，对全局内存的读取将首先尝试经过L1缓存；如果未命中，则接着尝试经过L2缓存；入股再次未命中，则直接哦才能够DRAM读取。一次数据传输处理的数据量在默认情况下是32字节。
关于全局内存的访问模式，有合并于非合并之分。
合并访问是指一个线程束对全局内存的一次访问请求(读或者写)导致最少次数的数据传输，否则称访问时非合并的。如果所有的数据传输中处理的数据都是线程束所需要的，那么合并度就是100%，即对应合并访问。
从全局内存转移到L2缓存的一片内存的首地址一定是一个最小粒度（这里是32字节）的整数倍。
使用cuda运行时API函数分配的内存的首地址至少是256字节的整数倍。
核函数常见的内存访问模式：
顺序的合并访问；
int n = threadIdx.x + blockIdx.x * blockDim.x;乱序的合并访问；
int tid_permuted = threadIdx.x ^ 0x1;(threadIdx^0x1是某种置换，作用是将0~31的整数做某种置换，交换两个相邻的数;)第一个线程块中的线程束将依然访问0~31个元素，只不过线程号于数组元素指标不完全一致。
不对齐的非合并访问；
int n = threadIdx.x + blockIdx.x * blockDim.x + 1;跨越式的非合并访问；
int n = blockIdx.x + threadIdx.x * gridDim.x;广播式的非合并访问；
在使用微软的编译器(MSVC)时有一个限制，即不能在核函数中使用在函数外部有const定义的浮点型常量。
从帕斯卡架构开始，如果编译器能够判断一个全局变量在整个核函数的范围都只可读，则会自动用函数__ldg()读取全局内存，从而对数据的读取进行缓存，缓解非合并访问带来的影响。而对于全局内存的写入，则没有类似的函数可用。所以，在不能满足读取和写入都是合并的情况下，一般来说，应当尽量做到合并地写入。
第8章 共享内存的合理使用
共享内存是一种可被程序员直接操控的缓存，其主要作用有两个：
一个是减少核函数中对全局内存的访问次数，实现高效的线程块内部通信；
另一个是提高全局内存访问的合并度；
8.1 例子：数组归约计算
在累加计算中出现了所谓的”大数吃小数“的现象。单精度浮点数只有6，7位精确的有效数字。
8.1.1 仅使用全局内存
对于多线程的程序，两个不同线程中指令的执行次序可能是和代码中所展现的次序不同。
__syncthreads();函数可保证一个线程块中所有（或者说所有线程束）在执行该语句后面的语句之前完全执行了该语句前面的语句。
在核函数中，位操作比相应的整数操作高效。
8.1.2 使用共享内存
所有的设备内存中，寄存器是最高效的，但在需要线程合作的问题中，用仅对单个线程可见的寄存器是不够的。需要使用对整个线程块可见的共享内存。
在核函数中，要将一个变量定义为共享内存变量，就要在定义语句中加上一个限定符__shared__， 一般情况下，我们需要的是一个长度等于线程块大小的数组。
在核函数中，定义一个共享内存变量，就相当于在每一个线程块中有了一个该变量的副本。
归约求和时候用共享内存相对于仅使用全局内存还有两个好处：
一是不在要求全局内存数组的长度N是线程块大小的整数倍；
另一个是在归约的过程中不会改变全局内存数组中数据。
8.1.3 使用动态共享内存
将前一个版本静态共享内存改成动态共享内存，只需要以下两个修改,但是不能指定数组大小。
（1）在调用核函数的执行配置中写入第三个参数<<<grid_size, block_size, sizeof(real) * block_size>>>；
(2) 还需要改变核函数中共享内存变量的声明方式 extern __shared__ s_y[]；
一般而言，在利用共享内存中的数据之前，都要进行线程块内的同步操作，以确保共享内存中的所有元素都已经更新完毕。
8.3 避免共享内存的bank冲突
关于共享内存，有一个内存bank的概念。为了获得高的内存带宽，共享内存在物理上被分为32个（刚好等于一个线程束中的线程数目，及内建变量warpSize的值）同样宽度、能被同时访问的内存bank。
只要同一个线程束内多个线程不同时访问同一个bank中不同层的数据，该线程束对共享内存的访问就只需要一次内存事务(memory transaction)。
当同一线程束内多个线程试图访问同一个bank中不同层的数据时，就会发生bank冲突。在一个线程束内对同一个bank中的n层数据同时访问将导致n次内存事务，称为发生了n路bank冲突。
同一个线程束中的32个线程(连续的32个threadIdx.x值)将对应共享内存数组S中跨度为32的数据。也就是说，这32个线程刚好访问同一个bank中的32个数据。通常可以用改变共享内存数组大小的方式来消除或减轻共享内存的bank冲突。
第9章 原子函数的合理使用
在cuda中，一个线程的原子操作可以在不受其他线程的任何操作的影响下完成对某个数据的一套“读-改-写”操作。该套操作也可以说是不可分的。
9.1 完全在GPU中进行归约
在归约求和的操作中，有两种办法可以在GPU得到最终的结果：
一个是用另一个核函数将较短数组进行下一步归约，得到最终结果；
另一个是原子操作；
9.2 原子函数
atomicAdd(address, val);第一个参数是待累加变量的地址，第二个参数是累加的值。 传给cudaMemcpy()函数的主机内存可以是栈内存。在传输少量数据时可以这样做，但在传输大量数据时这样做不安全，因为栈的大小是很有限的。
原子函数对它第一个参数指向的数据进行一次‘读-改-写’的原子操作，即一气呵成，不可分割操作。第一个参数可以指向全局内存，也可以指向共享内存。
原子函数没有同步功能。
从帕斯卡架构开始，在原来的原子函数基础上，引入两类新的原子函数。例如对于atomicAdd()来说，分别是
atomicAdd_system();和
atomicAdd_block();，前者将原子函数的作用范围扩展到整个同节点的异构系统(包括主机和所有设备)，后者将原子函数的作用范围缩小至一个线程块。
atomicAdd对数据类型支持是最全面的。atomicAdd()对双精度浮点数类型double的支持始于帕斯卡架构，对含有两个半精度浮点数变量的结构类型的支持也始于帕斯卡架构。
在所有的原子函数中，atomicCAS是比较特殊的，因为其他的原子函数都可以用它来实现。
使用原子函数时会引入随机性，因为原子函数只能保证每个原子操作的完整性，并不能保证不同的原子操作之间具有特定的次序。这种随机性会导致程序每次运行时结果都不一样，对测试造成了一定困难。
第10章 线程束基本函数与协作组
10.1 单指令-多线程执行模式
一个线程块不会被分配到不同的SM中，而总是在一个SM中，但一个SM可以有一个或多个线程块。不同的线程块之间可以并发或顺序执行，一般来说，不能同步(即使利用后面要介绍的协作组也只能在一些特殊的情况下进行线程块之间的同步)。
一个SM以32个线程为单位产生、管理、调度、执行线程。
在伏特架构之前，一个线程束中的线程拥有同一个程序计数器，但各自有不同的寄存器状态，从而可以根据程序的逻辑判断选择不同的分支。虽然可以选择分支，但是在执行时，各个分支是依次顺序执行的。在同一个时刻，一个线程束中的线程只能执行一个共同的指令或者闲置，这称为SIMT(单指令多线程)。
当一个线程束中的线程顺序地执行判断语句中地不同分支时，我们称发生了分支发散。分支发散是针对同一个线程束内部的线程的。
符合判断的线程执行，不符合判断的线程闲置。所以编写核函数时要尽量避免分支发散。
从伏特架构开始，引入了独立线程调度(independent thread scheduling）机制。每个线程有自己的程序计数器，有了新的线程束同步与通信模式。提高了编程的灵活性，降低了移植已有CPU代码的难度。缺点是增加了寄存器负担：单个线程的程序计数器一般需要两个寄存器。伏特架构的独立线程调度机制使得SM中每个线程可利用的寄存器少了两个。另外，独立线程调度机制使得假设了线程束同步(warp synchronous)的代码变得不安全。
也就是说__syncthreads()便不在是安全的做法了。__syncwarp()的粒度更细。
如果在伏特或者更高架构的GPU使用了线程束同步的程序，则可以在编译时候将虚拟架构指定为低于伏特架构的计算能力。
10.2 线程束内的线程同步函数
__syncwarp()束内同步函数，其原型是void syncwarp(unsigned mask = 0xffffffff)。默认值的全部32个二进制位都为1，代表线程束中的所有线程都参与同步。如果要排除一些线程，可以用一个对应的二进制位0的掩码参数。
当offset>=32时，我们在每一次折半求和后使用的线程块同步函数__syncthreads()；当offset<=16时，我们在每一次折半求和后使用束内同步函数__syncwarp()。
其中需要加上if (tid < offset)来避免16号线程同时被写和读，造成竞争情形。
10.3 更多的线程束内的基本函数
线程束表决函数(warp vote functions)， 线程束匹配函数(warp match functions)， 线程束洗牌函数(warp shuffle functions)及线程束矩阵函数(warp matrix functions)。
利用线程束洗牌函数进行归约计算 。其中有两个不同。第一，在进行线程束内的循环之前，将共享内存中的数据复制到了寄存器。第二用y += __shfl_down_sync(FULL_MASK, y, offset)取代了代码块，也就是说，去掉了同步函数，也去掉了对线程号的限制，因为洗牌函数能够自动处理同步与读-写竞争问题。对全部参与的线程来说，上述洗牌函数总是先读取各个线程中y的值，在将洗牌操作的结果写入各个线程的y。
10.4 协作组
协作组(cooperative groups)可以看作线程块和线程束同步机制的推广，它提供了更为灵活的线程协作方式，包括线程块内部的同步与协作、线程块之间(网格级的)同步与协作及设备之间的同步与协作。
协作组在cuda9.0才被引入，但关于线程块级的协作组功能可以用于开普勒架构及以上，而其他级别的协作组功能需要帕斯卡及以上架构才能通用。
协作组最基本的类型时线程组thread_group。thread_block的导出类型。
可以用tiled_partition()将一个线程块划分位若干片tile。
10.5 数组归约程序的进一步优化
当我们使用大小位128的线程块时候，当offset等于64后，只用了1/2的线程进行计算。故归约过程中线程平均利用率很低。如果能够提高归约之前所作计算的比例，则应该可以从整体上提升对线程的利用率。为了提高归约之前所作计算的比例，我们可以在归约之前将多个全局内存数组的数据累加到一个共享内存的数组的一个元素中。
这里要注意的时，千万不要让一个线程处理相邻的若干数据，因为这必将导致全局内存的非合并访问。要保证对全局内存的合并访问，则同一个线程所访问的数据之间必然具有某种跨度。该跨度可以是一个线程块的线程数，也可以是整个网络的线程数。
10.5.2 避免反复分配和释放设备内存
设备内存的分配与释放比较耗时。一种优化方案是使用静态全局内存代替这里的动态全局内存，因为静态内存是编译器间就会分配好的，不会在运行程序时反复地分配，故比动态内存分配高效很多。则可以在函数外部定义所需地静态全局内存变量。 __device__ real static_y[GRID_SIZE]，然后再用cudaGetSymbolAddress()获得一个指向该静态全局内存的指针，供核函数使用。
尽量避免在较内层循环反复地分配和释放设备内存。
第11章 CUDA流
cuda程序的并行层次主要有两个，一个是核函数内部的并行，一个是核函数外部的并行。
核函数外部的并行主要是指：
核函数计算与数据传输之间的并行；
主机计算与数据传输之间的并行；
不同的数据传输之间的并行；
核函数计算与主机计算之间的并行；
不同核函数之间的并行；
要获得较高加速比，需要尽量减少主机与设备之间的数据传输及主机中的计算，尽量在设备中完成所有计算，那么前四种核函数外部并行则不那么重要了。另外，若单个核函数的并行规模以及足够大了，在同一个设备中同时运行多个核函数也不会带来太多性能提升。
11.1 CUDA流概述
一个cuda流指的是有主机发出的在一个设备中执行的cuda操作序列。除主机端发出流外，还有设备端发出的流。
一个cuda流中各个操作的次序是由主机控制的，按照主机发布的次序执行。然而，来自于两个不同cuda流中操作不一定按照某个次序执行，而有可能并发或交错的执行。
任何cuda操作都存在于某个cuda流中，要么是默认流(default stream)，也称空流(null stream), 要么明确指定的非空流。
cudaStreamCreat(cudaStream_t*)创造流，cudaStreamDestory(cudaStream_t)销毁流。
为了实现不同cuda流之间的并发，主机在向某个CUDA流中发布一系列命令之后必须马上获得程序的控制权，不用等待cuda流中的命令在设备中执行完毕。
为了检查一个cuda流中所有操作是否都在设备中执行完毕，cuda运行时API提供了两个函数。
cudaStreamSynchronize()会阻塞主机，直到cuda流stream中所有的操作都执行完毕。cudaStreamQuery()不会阻塞主机，只是检查cuda流stream中的所有操作是否都执行完毕。
另外，也可考虑多个主机线程操控多个cuda流。
11.2 在默认流中重叠主机和设备计算
不同的是，核函数的启动时异步的(asynchronous)，或者说是非阻塞的(non-blocking)，意思是主机发出命令之后，不会等待该命令执行完毕，而会立刻获得程序的控制权。
虽然一个默认流就可以实现主机计算和设备计算的并行，但是要实现多个核函数之间的并行必须使用多个Cuda流。这是因为cuda流中的cuda操作在设备中是顺序执行的，故同一个cuda流中的核函数也必须在设备中顺序执行。
my_kernel<<<N_grid, N_block, N_shared, stream_id>>>(arguments)
利用Cuda流并发多个核函数可以提升GPU硬件的利用率，减少闲置的SM。
11.3 用非默认CUDA流重叠多个核函数的执行
要实现核函数执行于数据传输的并发，必须让这两个操作处于不同的非默认流，而且传输必须使用cudaMemcpy()函数的异步版本，即cudaMemcpyAsync()函数。
11.4 用非默认CUDA流重叠核函数的执行与数据传递
在使用异步的数据传输函数时，需要将主机内存定义位不可分页内存(non-pageable memory)或者固定内存(pinned memory)。
操作系统有权在一个程序运行期间改变程序中使用的可分页主机内存的物理地址。而不能修改不可分页内存的物理地址。
如果将可分页内存传输给cudaMemcpyAsyc()函数，则会导致同步传输。主机内存为可分页内存时，数据传输过程在使用GPU中的DMA之前，必须将数据从可分页内存移动到不可分页内存，从而必须于主机同步。主机无法在发出数据传输的命令后立刻获得程序的控制权，从而无法是不同的cuda流之间的并发。
不可分页内存的分配函数：
cudaMallocHost(void** ptr, size_t size):
cudaHostAlloc(void** ptr, size_t size, size_t flags)
书本例子多流并行没有得到最高的加速比的原因主要有两个：
Cuda操作的执行时间并不完全一样；
每个cuda流的第一个操作都是主机向设备传输数据，它们无法并发的执行。
第12章 使用统一内存
统一内存(unified memory)编程模型由cuda 6引入，从开普勒架构开始就用。
12.1 统一内存简介
12.1.1 统一内存的基本概念
统一内存是一种逻辑上的概念，它既不是显存，也不是主机的内存，而是一种系统中的任何处理器(CPU OR GPU)都可以访问，并能保证一致性的虚拟存储器。
统一内存编程从帕斯卡架构之后，功能加强很多。主要时因为这些架构的GPU具有了精细的缺页异常处理(page-fault handling)能力。
12.1.2 使用统一内存对硬件的要求
使用统一内存的硬件要求：
对于所有的功能，GPU的架构不低于开普勒架构，主机应用程序必须为64位的；
Windows系统都只能用第一代统一内存的功能。Linux系统都能用；
在具有IBM Power 9 和NVLink的系统中，伏特架构的GPU支持设备访问任何主机内存，包括用malloc分配的动态数组和栈上分配的静态数组。
12.1.3 统一内存编程的优势
使cuda编程更简单；
可能会比手工移动数据更好的性能。底层的统一内存实现，可能会自动将一部分数据放置到离某个存储器更近的位置；
允许GPU在使用统一内存的情况下，进行超量分配。仅限Linux系统和帕斯卡架构以及上；
12.2 统一内存的基本使用方法
统一内存设备中是当作全局内存使用的，而且必须在主机端定义或分配内存，而不能在设备端(核函数和__device__ （）函数)定义或分配。
12.2.1 动态统一内存
cudaMallocManaged(void** devPtr, size_t size, unsigned flags = 0), 该参数默认值是cudaMemAttachGlobal,表示分配的全局内存可由任何设备通过任何cuda流访问。只能在主机端使用该函数分配统一内存，而不能在核函数中使用该函数。
同一个程序能同时使用统一内存和非统一内存。
第一代统一内存来说，主机和设备不能并发地访问统一内存。又因为核函数地调用时异步地，故在调用任何核函数之后，不能紧接着从CPU访问任何统一内存变量(不管核函数是否访问了CPU将要访问的数据)，必须在中间加上一个同步函数。
12.2.2 静态统一内存
要定义静态统一内存，只需要在修饰符 __device__ 的基础上在加上__managed__修饰符即可。
仅仅在CPU中访问统一内存的话，在使用完主机内存后不会自动使用设备内存。所以，若将主机和设备的内存都纳入统一内存免责需要及时地让GPU访问统一内存或者用cudaMemPrefetchAsync()函数实现数据从主机到设备的迁移。
为了在使用统一内存时获得较高性能，需要避免缺页异常，保持数据的局部性（让相关的数据尽量靠近相应的处理器）但避免内存抖动（即频繁地在不同地处理器之间传输数据），cudaNenAdvise() or cudaMemPrefetchAsync()。
一般来说，在使用统一内存时，要尽可能多地使用cudaMemPrefetchAsync()函数，将缺页异常次数最小化。
