CUDA Programming Guide 学习笔记

2021-08-08 16:02:50

CUDA学习笔记

GPU架构

GPU围绕流式多处理器(SM)的可扩展阵列搭建，每个GPU有多个SM，每个SM支持数百个线程并发执行。目前Nvidia推出了6种GPU架构（按时间顺序，详见下图）：Fermi、Kepler、Maxwell、Pascal、Volta和Turing，每种架构的SM构成不尽相同。

CUDA Programming Guide 学习笔记

以Fermi架构为例，SM通常包含以下组成部分：

  • 32个CUDA核心，每核心由一整数算术逻辑单元ALU和一浮点运算单元FPU组成
  • 共享内存/L1Cache
  • 寄存器
  • 加载/存储单元
  • 特殊功能单元SFU（正弦余弦平方根等）
  • 两线程束调度器和两指令调度单元

CUDA采用单指令多线程（SIMT，Single Instruction Multiple Thread）架构来管理和执行线程，这些线程以32个为单位组成一个线程束（warps）。warp中所有线程并行执行相同的指令，但同一个warp中的线程可以以任意顺序执行，原因在于每个线程都拥有它自己的指令地址计数器和状态寄存器，因此尽管是在同一个warp中，每个线程也可以拥有自己独立的执行路径。

由于寄存器和共享内存的限制，SM中的active warps是有限的，每个线程消耗的寄存器或共享内存越多，则可以放在一个SM中的线程束就越少，极端的，若每个SM没有足够的寄存器或共享内存去处理至少一个块，那么内核将无法启动。当一个warp空闲时，SM就可以调度驻留在该SM中的另一个可用warp。由于线程束的本地执行上下文（计数器、寄存器、共享内存）都已被分配到所有的线程和线程块，在整个线程束的生存期中都是保存在芯片内的。当需要上下文切换时，所需要的操作就是将指向当前寄存器的选择器（或指针）更新，以指向下一个执行的线程束的寄存器，因此在并发的warp之间切换是没什么消耗的。

CUDA的线程管理

由一个kernel启动所产生的所有线程统称为一个网格(grid)，同一grid中的所有线程共享全局内存。

一个grid由多个线程块（block）构成，一个block包含一组线程（thread）。变量gridDim表示grid的大小，即一个grid中有多少block，而blockDim表示block的大小，即一个block中有多少thread。blockIdx用来获取block在grid中的索引，threadIdx用来获取thread在block中的索引。这些变量都是dim3类型的，是一个包含3个无符号整数（x，y，z）的结构体。

当一个kernel启动后，所产生的所有线程会被分配到若干个SM中执行，一个SM可以同时拥有多个block，但一个block只会由一个SM调度，block被分配给SM的顺序由块ID来确定，block中的所有线程都被分成了线程束。一个block一旦被分配好SM，该block就会一直驻留在该SM中。一旦所有的SM都被完全占用，所有剩余的线程块都保持不变直到当前的执行被完成，一旦一个线程块执行结束，将为该SM分配另一个线程块。同样的，block中的线程并不是物理上的同时执行，每个线程可以有不同的步调，CUDA提供了__synchreads()函数确保同一block内的所有线程保持同步，但不同block之间没有线程同步，不同block的线程之间也不允许相互同步， 前一个kernel死亡是唯一的全局同步方式的原因 ，GPU可以以任意顺序执行block。可以调用cudaDeviceSynchronize()强制CPU等待所有的核函数执行完成。

虽然同为block，但从逻辑角度看来，block是线程的集合，可以被组织为一维、二维或三维布局，每block最多1024个线程。而从硬件角度来看，block是一堆线程束的集合，呈一维布局，每32个连续线程组成一线程束。不同计算能力的设备常驻线程数目不同，详情见下表：

计算能力 3.0 3.2 3.5 3.7 5.0 5.2 5.3 6.0 6.1 6.2 7.0 7.5
常驻block数/每SM 16 16 16 16 32 32 32 32 32 32 32 16
常驻warps数/每SM 64 64 64 64 64 64 64 64 64 64 64 32
常驻thread数/每SM 2048 2048 2048 2048 2048 2048 2048 2048 2048 2048 2048 1024

当同一线程束中的线程执行不同的指令时，我们称这种现象为线程束分化，此时线程束将连续执行每一个分支路径，而禁用不执行这一路径的线程。线程束分化会大大削弱并行的性能，因此应避免在同一线程束中有不同的执行路径，我们可以通过线程束方法即\((tid/WARP\_SIZE)\%2==0\)而非直接通过线程判断即\(tid\%2==0\)。相同线程束中的线程可由连续的threadIdx.x来确定。

一般这样设置线程块：假如矩阵维度为\(Mat_{row\times col}\)，则线程大小可设置为：

dim3 grid((row + WARPS_IN_BLOCK - 1) / WARPS_IN_BLOCK);

dim3 block(WARP_SIZE, WARPS_IN_BLOCK);

这样设计的一个目的在于保证全局内存事务的对齐访问。对于这个grid而言，其grid维度为：

gridDim.x = (row + WARPS_IN_BLOCK - 1) / WARPS_IN_BLOCK;

gridDim.y = 1;

gridDim.z = 1;

其block维度为：

blockDim.x = WARP_SIZE;

blockDim.y = WARPS_IN_BLOCK;

blockDim.z = 1;

在核函数内部，矩阵的维度通过以下代码获取：

int row = blockIdx.x * blockDim.y + threadIdx.y;

int col = threadIdx.x + index * WARP_SIZE;

尽管CPU和GPU的变量在同一个文件作用域中被声明，但主机端代码不能直接访问设备变量（需通过CUDA运行时API），设备代码（如CUDA库函数、核函数）也不能直接访问主机变量特别是对于取地址&操作，切忌在主机端的设备变量中使用&运算符，因为它只是一个在GPU上表示物理位置的符号，并不是GPU内存上的地址。当然CUDA运行时API提供了一个函数用来获取一个设备全局变量（仅支持全局内存上的变量！）的地址：

cudaError_t cudaGetSymbolAddress(void** devPtr,const void* symbol);

CUDA内存管理

CUDA内存模型

一个核函数中的线程都有自己私有的本地内存和寄存器。一个线程块有自己的共享内存，对同一线程块中所有线程都可见，其内容持续线程块的生命周期。所有线程都可以读写全局内存，但只能读取常量内存和纹理内存，其内容持续整个程序的生命周期。

寄存器是GPU上运行速度最快的内存空间，寄存器对于每个线程而言都是私有的，与核函数的生命周期相同。若一个核函数使用了超过硬件限制数量的寄存器，则会用一级缓存或本地内存（Fermi架构）替代多余的寄存器。

  • 核函数中声明的一个没有其它修饰符的自变量通常存储在寄存器中，可以利用register 修饰符显式声明
  • 核函数中声明的索引引用为常量且能在编译期确定的数组也存储在寄存器中

不同计算能力的设备的寄存器（仅限32位的寄存器）数目不同，详情见下表：

计算能力 3.0 3.2 3.5 3.7 5.0 5.2 5.3 6.0 6.1 6.2 7.0 7.5
KB/每SM 64 64 64 128 64 64 64 64 64 64 64 64
KB/每Block 64 32 64 64 64 64 32 64 64 32 64 64
/每thread 63 255 255 255 255 255 255 255 255 255 255 255

本地内存实际上与全局内存在同一块存储区域，其用来存储索引引用未知的的本地数组、较大的本地结构体或数组、寄存器溢出的变量。寄存器和本地内存具有很多共同点，它们都是存储静态变量，也就是编译期能确定大小的变量；其次它们都是隐式声明的，即声明前不需要加任何修饰符。

共享内存，核函数中通过__shared__修饰符修饰的变量都被分配到共享内存上，其生命周期持续整个block，当一个block执行结束后，其分配的共享内存将被释放并重新分配给其它block。共享内存是块内线程间通信的基本方式，通过__synchreads()函数同步。每个SM都有一定数量的由线程块分配的共享内存，它具有高带宽低延迟的特点。共享内存可以通过固定大小创建（支持一维、二维和三维数组），也可以通过动态申请创建（仅限于未定大小的一维数组），但后者要求在声明时需要加上extern前缀extern __shared__ int num[];（核函数内），且在调用内核时需要在<<<>>>内加上第三个参数来指明需分配的共享内存的字节大小，两种声明均可在某个核函数内部或所有核函数外部进行。不同计算能力的设备的共享内存大小不同，详情见下表：

计算能力 3.0 3.2 3.5 3.7 5.0 5.2 5.3 6.0 6.1 6.2 7.0 7.5
KB/每SM 48 48 48 112 64 96 64 64 96 64 96 64
KB/每Block 48 48 48 48 48 48 48 48 48 48 96 64

常量内存，在全局空间内和所有核函数之外通过__constant__修饰符声明的变量存储在常量内存中，对所有计算能力的设备，都只可声明64KB的常量内存（因为常量内存是通过16位地址进行访问的，而16位地址能够快速进行访问），其驻留全局内存中，并在每个SM专用的常量缓存中缓存。常量内存是静态声明的，其生存期与应用程序的生存期相同，对网格内的所有线程都是可读的，且可以通过运行时函数cudaMemcpyToSymbol对主机可写，该函数声明如下：

cudaError_t cudaMemcpyToSymbol(const void* symbol,const void* src,size_t count);

该函数将count个字节从src指向的内存复制到symbol指向的常量内存（或全局内存，函数不会对这片内存区域作区分）中。由于每从一个常量内存中读取一次数据，都会广播给线程束里的所有线程，因此对于数学公式中的系数等需要所有线程从相同内存地址读取的数据而言，常量内存表现很好。因此，对常量内存而言，若线程束中的所有线程都访问相同的位置，则这个访问模式就是最优的，相反若线程束中的线程访问不同的地址，则访问需要串行，一个常量内存的读取成本取决于线程束中线程读取相同地址的数目。

Tip:如果一个常量只是字面值，那么最好用 #define 对字面值进行定义，因为这样可以减少常量内存的使用。

纹理内存同样驻留在全局内存中，并在每个SM的纹理缓存中缓存。纹理缓存是对二维空间局部性的优化，通过纹理缓存访问二维矩阵的邻域会获得加速。纹理可以是一段连续的全局内存，也可以是一个CUDA数组，但是CUDA数组对局部寻址有优化，称为“块线性”，原理是将邻域元素缓存在同一条cache线上，这将加快邻域内的寻址，缺点在于将数据拷贝至CUDA数组非常耗时。只读缓存可以通过__ldg函数和const __restrict__ *限定指针来向编译器指出该数据在内核中是只读的。只读缓存适用于分散读取。

全局内存是GPU中最大、延迟最高且最常用的内存，生命周期持续整个应用程序，其声明可以在任何SM设备上被访问到，当然也就可以访问所有核函数中的所有线程（注意block间无法同步）。一个全局内存变量可以通过__device__修饰符静态声明，此类变量通过cudaMemcpyToSymbol初始化并通过cudaMemcpyFromSymbol取回，也可以通过cudaMalloc函数动态分配（当然得记得cudaFree释放）。

GPU缓存包括一级缓存、二级缓存、只读常量缓存、只读纹理缓存，它们都是不可编程的缓存。每个SM都有一个一级缓存，所有SM共享一个二级缓存，两者都被用来存储本地内存和全局内存中的数据，包括寄存器溢出的部分。但要注意的是，GPU上只有内存加载操作可以被缓存，内存存储操作不能被缓存。一级缓存和共享内存共享同一硬件资源，可通过cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig)或cudaError_t cudaFuncSetCacheConfig(const void *func, enum cacheConfig)按核函数的需求进行配置。不同计算能力的设备的常量缓存、纹理缓存大小不同，详情见下表：

计算能力 3.0 3.2 3.5 3.7 5.0 5.2 5.3 6.0 6.1 6.2 7.0 7.5
常量缓存KB/每SM 8 8 8 8 8 8 8 4 8 8 8 8
纹理缓存KB/每Block 12~48 12~48 12~48 12~48 12~48 12~48 12~48 24~48 24~48 24~48 32~128 32 or 64
声明 存储位置 作用域 生命周期
float var 寄存器 线程 线程
float var[100] 本地内存 线程 线程
__shared__ float var 共享内存 block block
__device__ float var或cudaMalloc((float**)&var,100) 全局内存 全局 程序
__constant__ float var 常量内存 全局 程序

固定内存并不是CUDA内存模型的一部分，甚至不在GPU上，固定内存存在的意义在于GPU不能在可分页（默认，malloc分配的就是可分页内存）主机内存上安全的访问数据（GPU无法控制主机在物理位置移动该数据）。当从可分页主机内存传输数据到设备内存时，CUDA首先分配固定的主机内存（通过页面锁定方式），将主机源的数据复制到固定内存中，然后从固定内存传输数据给设备内存。cudaMallocHost函数允许在代码中显式分配固定主机内存，函数声明如下：

cudaError_t cudaMallocHost(void** devPtr,size_t count);

cudaError_t cudaFreeHost(void* ptr);

需要特别注意的是，尽管固定内存在主机上，但我们仍不能在主机端代码中直接访问！相比直接通过可分页内存读取，固定内存的分配和释放成本更高，但它能用比前者高得多的带宽进行读写，因此适用于大规模数据传输。但由于固定内存占用了主机虚存的可分页内存数量，因此分配过多的固定内存会降低主机系统的性能。

零拷贝内存与固定内存类似，同样不是CUDA内存模型的一部分，实际上，它就是固定内存，但该内存映射到了设备地址空间中。主机和设备都可以访问零拷贝内存！当然，这也就意味着，你必须时刻同步主机和设备间的内存访问，而由于每一次读写操作都需要进过PCIe总线进行传输，因此零拷贝内存不适应于频繁的数据读写。零拷贝内存通常作为设备内存不足时的补充。零拷贝内存通过cudaHostAlloc函数创建，并通过cudaHostGetDevicePointer获取映射后的设备指针，函数声明如下：

cudaError_t cudaHostAlloc(void** hostPtr,size_t count,unsigned int flags);

cudaError_t cudaHostGetDevicePointer(void** devPtr,void* hostPtr,unsigned int dFlag);

cudaError_t cudaFreeHost(void* ptr);

对于零拷贝内存而言，这儿的flags标志为cudaHostAllocMapped。注意cudaHostGetDevicePointer函数的最后一个参数dFlag必须被设置为0。通过主机端指针hostPtr和设备端指针devPtr，CPU和GPU可以访问位于主机内存上的同一片区域。由于每次读写都存在较大延迟，不建议使用。

统一虚拟寻址（UVA）是一种支持主机内存和设备内存共享同一个虚拟地址空间的寻址方式。这也就意味着，通过UVA，由cudaHostAlloc分配的固定主机内存具有相同的主机和设备指针，因此可以将返回的指针直接传递给核函数。需要注意的是，UVA只是提供了一个单一的虚拟内存地址空间，它不会将数据从一个物理位置转移到另一个位置。事实上，UVA只是免去了上述零拷贝内存中的cudaHostGetDevicePointer的步骤而已，数据仍然是在主机固定内存上的，同样不建议使用。

统一内存寻址相比UVA，创建了一个托管内存池，内存池中已分配的空间可以用相同的指针在CPU和GPU上进行访问，主机和设备之间的数据传输将由统一内存自动管理。托管内存在主机内存和设备内存上同时分配空间，并通过UVA使得两者共享同一指针。相比未使用统一内存寻址技术的程序，统一内存寻址消除了重复指针（无需区分主机端与设备端）、提供了自动数据传输功能，换句话说，后者比前者少了手动copyHost2Device和copyDevice2Host操作，实际上，这些操作被推迟到了发生页面故障时。使用托管内存并不会造成多少性能损失（实际上，托管内存的初始化时间很长，但核函数和数据传输速度几乎一样），但可以大大降低编程工作量。托管内存同样既可以通过__managed__修饰符静态分配（只支持__device__全局变量），也支持cudaMallocManaged动态分配，函数声明如下：

cudaError_t cudaMallocManaged(void** ptr,size_t size,unsigned int flags=0);

这儿的指针ptr在主机端代码和设备端代码中都是有效的，且行为和功能与未使用托管内存的指针一致。若系统拥有多个GPU设备，则托管应用需要附加的步骤，因为托管内存分配对系统中的所有设备都是可见的，需要设备环境变量export CUDA_VISIBLE_DEVICES=0（设定哪个设备对程序可见）来使得托管内存只分配在一个设备上。

全局内存

CUDA中的指令都是以线程束为单位进行发布和执行的，内存存取指令也一样。线程束中的每个线程都提供了一个加载或存储的内存访问请求，根据线程束中内存请求地址的分布，内存访问分为若干种模式。

全局内存是通过缓存来实现加载/存储的，所有对全局内存的访问都会首先尝试通过一级缓存，若一级缓存缺失，该请求转向二级缓存，若二级缓存仍然缺失，则由内存DRAM完成。

对于一级缓存而言，内存访问是由一个128字节的内存事务实现的，换句话说，对于每次内存访问请求，一级缓存都将从设备内存获取一个128字节的对齐段。若内存访问的第一个请求地址是128的倍数，则此时内存访问是对齐的，若线程束中的这些线程访问的是一个连续的内存块时，则此时内存访问将合并。非对齐或未合并的内存访问会造成带宽浪费，此时内存事务获取的部分字节都将不能被使用。不同于CPU的一级缓存，GPU一级缓存专为空间局部性设计，并未考虑时间局部性，这也就意味着，频繁访问并不会增加该数据留在缓存中的概率。

相比128字节的一级缓存，仅使用二级缓存可以提供更小的只有32字节的内存事务，因此对于非对齐或未合并的内存访问请求而言，后者有更高的总线利用率（即使是在最差情况下）。然而，实验表明仅使用二级缓存并不能减少缓存的整体价值时间，若使用一级缓存，一个非对齐访问的一次内存事务请求会将多余数据存到一级缓存，这些数据可以用于后续的非对齐内存访问。若仅使用二级缓存，那么每一次非对齐请求需要多个内存事务，且对之后的请求没有作用。

相比内存的读取操作，内存的存储操作只能通过二级缓存，因此存储操作是通过32字节的内存事务实现的。但内存事务可以同时被分为一段、两段和四段，这样做的目的在于，若两内存存储地址属于同一128字节区域，但不属于同一对齐的64字节区域，此时执行一个四段事务（一个128字节操作）比执行两个一段事务（两个32字节操作）效果更好。因此内存存储操作同样需要合并对齐。

TIP:相比数组结构体（AoS，数组每个元素是一个结构体），结构体数组（SoA，结构体每个成员是数组）能更充分的利用GPU的内存。原因在于SoA被相同内存操作引用的同字段数据元素在存储时是彼此相邻的，加载和存储内存请求不会重复，每次访问都由一个独立的内存事务来处理。

在分析内存访问性能时，有两个指标需要注意：内存访问，即完成一次独立内存请求的时间；内存带宽，即SM访问设备内存的速度，以每单位时间内的字节数计。优化设备内存带宽利用率有两个目标：

  • 对齐及合并内存访问，以减少带宽的浪费
  • 足够的并发内存操作，以隐藏内存延迟

后者可以通过循环展开使得每个线程产生更多的独立内存访问，也可以通过修改核函数启动的执行配置使得每个SM有更多的并行性。

共享内存

共享内存是片上内存，和一级缓存在物理上更接近SM，因此具有比全局内存更低的延迟和更高的带宽。通常，我们可以将之视为一个可编程管理的缓存，当数据移动到共享内存中以及数据被释放时，我们对其有充分的控制权。物理上，每个SM都有一个小的低延迟内存池，这个内存池被当前正在该SM上执行的线程块中的所有线程所共享。

共享内存使同一线程块中的线程能够相互协作。当每个线程块开始执行时，会分配给它一定数量的共享内存，这个共享内存的地址空间被线程块中所有的线程共享，生命周期与创建时所在的线程块相同。共享内存被SM中的所有常驻线程块划分，一个核函数使用的共享内存越多，处于并发活跃状态的线程块就越少。

当使用共享内存设计核函数时，应该重点注意两个概念：

  • 跨存储体映射元素
  • 从线程索引到共享内存偏移的映射

存储体冲突

共享内存是一个一维地址空间，它被分为32个（线程束大小）同样大小、可以被同时访问的内存模型，称为存储体。存储体存在两种地址模式：32位和64位。以32位的Fermi设备为例，每个存储体在每两个时钟周期内都有32位的带宽，连续的32位字映射到连续的存储体中，存储体的索引可以通过\(存储体索引=(地址\div 4字节)\%32\)公式得到。地址模式可以通过cudaDeviceGetSharedMemConfig查询，pConfig的值为cudaSharedMemBankSizeFourByte或cudaSharedMemBankSizeEightByte。更改存储体大小不会增加共享内存的使用量，也不会影响核函数的占用率。

cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig *pConfig)

当线程束发出共享内存请求时，存在如下3种典型的访问模型：

  • 并行访问：多个地址访问多个存储体
  • 串行访问：多个地址访问同一个存储体
  • 广播访问：单一地址读取单一存储体

CUDA Programming Guide 学习笔记

当多个地址请求落在相同的内存存储体中时，就会发生存储体冲突，这会导致请求被重复执行，硬件会将存储体冲突的请求分割到尽可能多的独立的无冲突事务中。最佳情况下，当每个地址都位于一个单独的存储体中时，只需要一个内存事务即可完成。最差情况下，当线程束32个线程全部访问同一存储体中不同的内存地址时，需要32个内存事务才能完成。此外，广播访问发生在线程束多个线程全部访问同一存储体中相同地址时，此时被访问的字就会被广播到所有请求的线程中--尽管广播访问也只需要一个内存事务，但由于只读取一小部分字节，所有带宽利用率很低。对于写访问而言，当线程束多个线程写入同一地址时，这个字只能由其中一个线程写入，执行这个写入操作的线程是不确定的。

尽管一级缓存和共享内存位于相同的片上硬件上，但一级缓存是通过缓存行而非存储体进行访问的，且数据删除工作由硬件完成。

同步

GPU线程在不同内存中写入数据的顺序，不一定和这些数据在源代码中访问的顺序相同，一个线程的写入顺序对其它线程可见时，它可能和写操作被执行的实际顺序不一致。此外若指令之间相互独立，则线程从不同内存中读取数据的顺序和读指令在程序中出现的顺序也不一定相同。

CUDA中提供了障碍和内存栅栏两个方法来执行块内同步。障碍void __syncthreads()要求块中的线程必须等待直到所有线程都到达该点，且确保被这些线程访问的所有全局和共享内存对同一块中的所有线程可见（ 一个线程对变量值的修改，能够及时的被其他线程看到 ）。需要注意的是，调用__syncthreads时必须保证块内所有的线程都能执行到该位置，否则很可能会导致块中的线程无限期的等待对方。例如以下代码可能会导致块中线程无限期的等待对方，因为块中的所有线程没有达到相同的障碍点：

if (threadId % 2 == 0) {

  __syncthreads();

} else {

  __syncthreads();

}

内存栅栏不执行任何线程同步，它只保证栅栏前的任何内存写操作对栅栏后的其他线程都是可见的。CUDA提供了三个内存栅栏函数：block级内存栅栏void __threadfence_block()、网格级内存栅栏void __threadfence()和跨系统（包括主机和设备）级内存栅栏void __threadfence_system()，后两者都会挂起调用的线程，直到全局内存、固定主机内存和其它设备内存中的所有写操作对同网格或全部设备中的线程和主机线程可见。

线程束洗牌

线程束洗牌（shuffle）使线程束中的线程彼此之间可以直接交换数据，换句话说，只要两个线程在相同的线程束中，那么就允许这两个线程直接读取（只能读）另一个线程的寄存器，而无需通过共享内存或全局内存来进行。洗牌指令比共享内存拥有更低的延迟，且无需消耗额外的内存。

CUDA提供的洗牌指令支持整型变量和浮点型变量，包括int, unsigned int, long, unsigned long, long long, unsigned long long, float, double，若包含了头文件cuda_fp16.h，则还支持__half, __half2类型（仅限cuda9.0及以上）。

以int类型为例介绍cuda提供的四种形式的洗牌指令，参数mask用于限定哪些线程参与洗牌，直接设定0xffffffff允许所有线程即可。

int __shfl_sync(unsigned mask, int var, int srcLane, int width=warpSize);

__shfl_sync指令使线程束中的每个线程都可以直接从由srcLane确定的线程中获取var变量值，并返回。srcLane的含义取决于width值，width值可被设置为2~32之间2的指数(\(2,4,8,16,32\))，默认值为warpSize即32，srcLane的值可以大于width。束内索引通过threadIdx.x % width得到，特定线程ID为threadIdx.x&~(width-1)+srcLane ，这个每width个线程可以执行单独的洗牌操作。当线程束中所有线程执行__shfl_sync的srcLane的值相同时，该函数执行线程束广播操作。

int __shfl_up_sync(unsigned mask, int var, unsigned int delta, int width=warpSize);

int __shfl_down_sync(unsigned mask, int var, unsigned int delta, int width=warpSize);

__shfl_up_sync指令通过减去delta（右移delta个）来计算源束内索引，并返回var变量值。类似的，__shfl_down_sync指令通过增加delta（左移delta个）来计算源束内索引。需要注意的是，多余的delta个线程将保持不变。

int __shfl_xor_sync(unsigned mask, int var, int laneMask, int width=warpSize);

__shfl_xor_sync指令通过按位异或laneMask来计算源束内索引。

CUDA流

CUDA流是一系列异步的CUDA操作，这些操作安装主机代码确定的顺序在设备上执行。流封装这些操作，保持操作的顺序，允许操作在流中排队，并使它们在先前的所有操作之后执行，且可以查询排队操作的状态。在同一个CUDA流中的操作严格按序执行，而在不同CUDA流中的操作不受限制。流中操作的执行相对于主机总是异步的，CUDA运行时决定何时可以在设备上执行操作。

基本操作

所有的CUDA操作都在一个流中显式或隐式地运行，若没有显式地指定一个流，则内核启动和数据传输默认使用空流，若显式设定了一个CUDA流cudaStream_t，则内核启动和数据传输在该非空流中异步进行。非空流通过cudaStream_t修饰符声明，通过cudaStreamCreate创建，通过cudaStreamDestroy释放。

cudaError_t cudaStreamCreate(cudaStream_t *pStream)

cudaError_t cudaStreamDestroy(cudaStream_t pStream)

非空流所有的操作都不阻塞主机执行，相反，隐式声明的空流与主机同步，换句话说，添加到空流上的操作会导致主机的操作阻塞。另一方面，非空流被进一步划分为阻塞流和非阻塞流，两者都不会阻塞主机操作，但空流可以阻塞阻塞流中的操作，具体而言，任何发布到阻塞流中的操作都被挂起等待，直到空流中先前的操作执行结束才开始执行，cudaStreamCreate创建的流是阻塞流。而非阻塞流不会被空流阻塞，CUDA中通过cudaStreamCreateWithFlags函数并设置flags值为cudaStreamNonBlocking来创建非阻塞流：

cudaError_t cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags)

在核函数执行配置中设置第四个参数为非空流即可在该流中启动内核，需要注意的是，由于所有CUDA流操作都是异步的，因此返回错误的API调用并不一定是产生错误的那个调用。另外，当cudaStreamDestroy被调用时，若该流中仍有未完成的工作，函数将立即返回，当流中所有的工作都已完成时，与流相关的资源将被自动释放。

CUDA还提供了两个函数用于检查流：cudaStreamSynchronize强制阻塞主机，直到给定流中所有的操作都已经完成；cudaStreamQuery则不会阻塞主机，只会在所有操作完成时返回cudaSuccess，否则返回cudaErrorNotReady。

cudaError_t cudaStreamSynchronize(cudaStream_t stream)

cudaError_t cudaStreamQuery(cudaStream_t stream)

CUDA流还具有优先级，高优先级流的网格队列可以优先占有低优先级流已经执行的工作，流优先级不会影响数据传输操作，只对计算内核有影响。可以通过cudaStreamCreateWithPriority创建特定优先级的流。

cudaError_t cudaStreamCreateWithPriority(cudaStream_t *stream, unsigned int flags, int priority)

由于硬件资源（PCIe总线）的限制，数据传输操作虽然分布在不同的流中，但并没有并发执行。具有双工PCIe总线的设备可以重叠两个数据传输，但它们必须在不同的流中以及不同的方向上。

CUDA事件

CUDA事件本质上是CUDA流中的标记，它与该操作流特定点相关联，主要用来检查正在执行的流操作是否已经到达了给定点，用于同步流的执行、监控设备的进展。CUDA事件通过cudaEvent_t修饰符声明，通过 cudaEventCreate创建，通过cudaEventRecord排队进入CUDA流，通过cudaEventDestroy销毁：

cudaError_t cudaEventCreate(cudaEvent_t *event)

cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0)

cudaError_t cudaEventDestroy(cudaEvent_t event)

类似的，当cudaEventDestroy函数被调用时，若事件尚未起作用，调用会立即返回，当事件被标记完成时会自动释放与该事件相关的资源。CUDA同样提供了两个函数用于检查事件：cudaEventSynchronize和cudaEventQuery：

cudaError_t cudaEventSynchronize(cudaEvent_t event)

cudaError_t cudaEventQuery(cudaEvent_t event)

除这两函数外CUDA还提供了cudaStreamWaitEvent函数用于使指定流等待事件，该事件无需在该流中：

cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event)

cudaStreamWaitEvent函数可以用于指定流间的同步。CUDA还提供了计时函数cudaEventElapsedTime用于计算两事件间CUDA操作的运行时间：

cudaError_t cudaEventElapsedTime(float *ms, cudaEvent_t start, cudaEvent_t stop)

此函数返回事件start和事件stop之间的运行时间，以毫秒为单位（异步延迟，不准确），注意两事件不要求在同一个CUDA流中。

流回调

流回调函数是由应用程序提供的一个主机函数，并在流中使用cudaStreamAddCallback函数注册：

cudaError_t cudaStreamAddCallback(cudaStream_t stream, cudaStreamCallback_t callbackFunc, void *userData, unsigned int flags)

userData参数用于传递给回调函数，flags参数目前没有任何意义，必须设置为零。回调函数callbackFunc通过如下方式声明：

void CUDART_CB callbackFunc(cudaStream_t stream, cudaError_t status, void *userData)

在流中所有先前排队的操作完成后，回调函数才能在主机上执行，每使用cudaStreamAddCallback函数一次，只执行一个回调，并阻塞队列中排在后面的工作，直到回调函数完成。当它被CUDA运行时调用时，回调函数会通过调用它的流，并且会有错误代码来表明是否有CUDA错误的发生。在回调函数中不可以调用CUDA的API函数，也不可以执行同步。

CUDA指令

指令是处理器中的一个逻辑单元，对于CUDA而言，影响CUDA内核生成指令的3大因素为：浮点运算、内置和标准函数、原子操作。

内部函数

标准函数主要包含C/C++标准库中的数学运算，支持主机和设备的访问和操作。相比之下，CUDA内部函数只能对设备代码进行访问，但在编译时却会被分解为等价标准函数更少的指令，因此内部函数速度更快，但精度更低。在编译时，若一个函数是内部函数，则编译器会对该函数产生更积极的优化和更专业化的指令生成。

CUDA编译器中有两种方法可以控制指令级优化类型：编译器标志、内部或标准函数调用。后者例如内部函数__fdividef与运算符/相比，在执行浮点数除法时速度更快但数值精确度相对较低。前者例如–fmad=true选项可以全局启动FMAX指令来优化性能（乘法加a*b+c运算通过MAD指令可以得到更快的速度），--use_fast_math选项将用等价的内部函数替换应用程序中所有的标准函数。

可以通过nvcc的--ptx标志让编译器在并行线程执行（PTX）和指令集架构（ISA）中生成程序的中间表达式，而不是生成一个最终的可执行文件。通过这个中间表达式我们可以将内部函数的优化可视化。

原子操作

一条原子操作用来执行一个数学运算，此操作是一个独立不间断的操作，且没有其它线程的干扰。原子操作分为算术运算函数、按位运算函数和替换函数。

CAS运算符（原子级比较交换运算符）是原子操作中最重要的一个操作，通过这个操作我们可以生成自己的原子操作。CAS将内存地址、期望值和新值作为参数，若内存地址当前存储值与预期值相等，则新值存入目标位置，否则不做任何事情，但无论如何CAS都将返回当前内存地址中的值，若返回值等于预期值则CAS操作一定成功了。

以加法为例，我们可以通过CAS操作实现原子级加法操作如下：

__device__ int atomicAdd(int *num, int inc) {

	int expect=*num;

    int old=atomicCAS(num,expect,expect+inc);

    while(old!=expect){

        expect=old;

        old=atomicCAS(num,expect,expect+inc);

    }

    return old;

}

原子操作对于数据竞争而言非常有必要，但同样需要付出很高的代价：

  • 一个对全局或共享内存的原子操作将不通过缓存，这意味着每个原子操作都将通过I/O进行数据传输
  • 竞争冲突可能要求发生冲突的线程不断重试
  • 当线程在同一个线程束中时必须执行不同的指令，线程束执行是序列化的

由于某些原子操作不支持所有的数据类型，因此有些数据类型的原子操作需要自己定义，有一个办法是用一个变量中支持的类型存储浮点数的原始比特位。以float为例，这里我们通过unsigned int类型来实现float类型的原子加法操作：

__device__ float atomicAdd(float *num, float inc) {

    unsigned int *typed=(unsigned int*)num;

    float cur=*num;

    unsigned int expect=__float2uint_rn(cur);

    unsigned int aim=__float2uint_rn(cur+inc);

    int old=atomiccAS(typed,expect,aim);

    while(old!=expect){

        expect=old;

        aim=__float2uint_rn(__uint2float_rn(old)+inc);

        old=atomiccAS(typed,expect,aim);

    }

    return __uint2float_rn(old);

}
上一篇：js插件---bootstrap插件daterangepicker是什么


下一篇：Delphi中TStringlist.count一个奇怪问题及其解决