什么是CUDA

CUDA是nvida基于自家GPU提供的一套并行计算框架，通过CUDA可以很方便的进行编码利用GPU强大的并行处理能力完成对应任务计算。

基本概念

HOST（主机部分）：CPU以及CPU所使用的内存以及在CPU上执行的代码。

DEVICE（设备）：GPU以及GPU所使用的显存以及对应运行在GPU上的代码。

基本逻辑结构如下图

Kernel(核函数)：GPU上执行的函数被称为核函数。

Thread（线程）：GPU所执行的最小单元（一般由GPU上的一个处理核执行）。

Block（线程块）：多个Thread组成的一组线程，Block内的线程可以进行数据共享，Block之间无法进行数据共享，不同的block是并行执行的，并且不同的block执行上没有先后顺序。在一个Block中，一般最多创建不超过512个Thread。

Grid（线程格）：有多个Block组成，一般一个Grid内的线程数量不超过65535个（该Grid内所有Block线程之和），一般一个GPU设备可以生成2到3个Grid。

Wrap：逻辑上，所有的thread都是并行执行，但从硬件的角度看，并不是所有的thread都能够在同一时间执行，Wrap是GPU调度和执行thread的最基本单元（可理解为GPU硬件执行单元），一般情况下，一个wrap包含32个thread，这32个thread执行同一条处理指令（数据不同）。此外一个wrap中被执行的线程必然属于同一个Block，如果Block中的线程数目不是wrap大小的整数倍，在block被执行时，也会被凑够为wrap的整数倍（多余的线程只是状态被设置为非活动状态）。

Warp Divergence (Wrap分歧)：CPU拥有非常好的分支预测能力，如果预测正确，CPU只会由很小的消耗，和CPU相比，GPU的分支预测就非常差，同一个wrap中的thread同时执行相同的指令，如果thread遇到分支控制语句后，不同thread进入不同的分支，就会导致其余分支都被阻塞，十分影响性能，这类问题就被称为Warp Divergence，Warp Divergence只会发生在同一个wrap中。

SM（Streaming Multiprocessor）：GPU硬件上的一个概念，一个GPU有非常多的SM组成，一个SM内部，会有该SM对应的寄存器，指令cache，指令buffer，wrap调度器，计算core等部分。一般情况下一个Wrap会被分配到一个SM上执行。当然一个SM可以执行多个Wrap。

上图是SM的组成部分，其中绿色的Core是SP，也可以称为CUDA

core，CUDA core是一个单精度微处理器，DP Unit是双精度微处理器。一个SM里面一般包含32个DP Unit，64个CUDA core（SP）。此外SM换包括LD/ST单元，LD/ST主要用于内存操作，SM中另外要给组件SFU（Special function unit）部分，主要用于执行一些编译器内建的特殊函数，例如cos等。

SP（Streamingprocessor）：流处理器，GPU中专门用于数学（整形，浮点）计算的处理单元，多个SP与其他组件一起组成一个SM。SP相当于一个微型处理器，基本结构图如下所示：

GPU内存架构

GPU内存逻辑结构如下：

LocalMemory：本地内存，属于线程私有内存，只能有当前线程访问，其他线程不可访问，当线程内的寄存器不够使用时，会动用该部分内存。

sharedMemory：共享内存，可以被同一个block中的所有线程访问，不同的线程可以通过共享内存进行通信。一般一个block内会对应一块共享内存，该共享内存只能被block内的线程访问（Block中的线程会被分配给同一个SM执行）。

Global memory：全局内存（占据显卡内存的绝大部分），CPU和GPU都可以访问（CPU通过PCIe总线进行访问），所有线程都可以访问Global memory。

常量内存：特殊类型的全局内存，Grid内的所有线程都可以只读访问。

纹理内存：特殊类型的全局内存，Grid内的所有线程都可以只读访问。

线程组织方式

线程首先被组织为线程块（thread

block），多个线程块被组织成Grid。

线程块内可以进行数据的共享，以及执行过程的同步（一个线程块内的某个线程可以被挂起，直到其他线程执行到同样的位置）（执行流程的同步）。

在线程块内，每个线程都有一个线程ID，根据线程ID可以进行复杂寻址。

多个线程块组成线程Grid，kernel被映射到Grid进行执行，因此一个Kernel可以使用的线程数量非常巨大。

同一个Grid内不同Block的线程彼此之间不能通讯和同步。在Grid内的每个Block都有要唯一性的ID进行标识，根据Block ID可以进行复杂寻址。

硬件在调度时将block划分为固定数量的Wrap，同一个block对应的wrap会被调度到某一个SM执行，不同Block中的线程可能被同一个SM执行，也可能被不同的SM执行。

CUDA编码

CUDA中，编码上通过编写kernel函数来让GPU完成计算任务。

CUDA函数

CUDA中的函数分为三类。

1：只能在CPU侧调用和执行的函数，该类函数通过在函数前添加__host__限定词限制。

2：在主机侧被调用，在GPU设备侧被执行的函数，这类函数通过在函数声明前添加__global__限定词限制。对于这类函数，返回值只能时void，并且不支持递归，参数的大小也被限制在256个字节以内，该类函数内也不能声明静态变量。对于被__global__修饰的函数，在调用时，需要指明执行线程的配置。

3：只能在GPU设备侧被调用的函数，这类函数通过__device__修饰。被__device__修饰的函数无法获取函数地址。

CUDA变量

对于变量而言，通过不同的修饰限定词，也区分为以下几种

1：被__device__修饰的变量，该类变量驻留在GPU设备全局内存空间中，在应用的整个生命周期中有效，Grid内的所有线程都可以访问该变量，主机侧通过runtime库也可以访问该变量。

2：被__constant__修改的变量，该类变量驻留在GPU侧的常量内存空间中，在应用的整个声明周期有效，Grid内的所有线程都可以访问该变量，主机侧通过runtime库也可以访问该变量。

3：被__shared__修饰的变量，该类变量驻留在Block线程块的共享内存中，在Block声明周期内有效，只有同Block内的线程可以访问。

CUDA函数的执行

通过__global__修饰的kernel函数，在主机侧被调用时，需要指定该函数如何在GPU侧被执行。

例如对于如下一个kernel函数

__global__ void

Func(float *param)，在调用的时候形式如下：

Func<<<Dg,

Db, Ns, S>>>(param);

其中<<<Dg, Db, Ns, S >>>部分描述了GPU上执行的线程组织方式。

Dg是一个dim3（x,y,z）的数据类型或int型，用于指定Grid的维数和大小。如果时int型，表示是一个一维的组织结构。

Db是一个dim3的数据类型或int型，用于指定Block的维数和大小。如果是int型，表示是一个一维的组织结构。

Ns是一个size_t类型的数据，用于指定每个block中可动态分配的共享内存的字节数量，默认为0.

S：cuda流类型，默认为0.

如下的调用方式

dim3 grid(3,2,1), block(4,3,1);

kernel_name<<<grid,

block>>>(...);

表示一个Grid中有3*2*1个Block，在（x,y,z）三个方向上排列方式分表为3/2/1;一个Block中有4*3*1个线程，在(x,y,z)三个方向上排列方式为4/3/1.

如下的kernel调用方式：

kernel_name<<<5,8>>>(...);

表示一个Grid中有5个Block，在(x,y,z)三个方向上的分布方式为5/1/1;在一个Block中有8个线程，在(x,y,z)三个方向上的分布方式为8/1/1。

在kernel函数中，可以通过内置变量计算线程ID

threadIdx.[x, y, z]表示Block内Thread的编号

blockIdx.[x, y, z]表示Gird内Block的编号

blockDim.[x, y, z]表示Block的维度，也就是Block中每个方向上的Thread的数目

gridDim.[x, y, z]表示Gird的维度，也就是Grid中每个方向上Block的数目

对于kernel_name<<<4, 8>>>(...)这样一维的结构，如下所示：

线程ID = blockIdx.x *

blockDim.x + threadIdx.x

对于如下二维结构

dim grid(4,1,1),

block(2,2,1);

kernel_name<<<grid,

block>>>(...)

线程ID计算如下：

int blockId =

blockIdx.x + blockId.y * gridDim.x;

int threadId =

blockId * (blockDim.x * blockDim.y) + (threadIdx.y *blockDim.x) + threadIdx.x;

对于如下多维结构kernel函数，线程ID计算方式如下：

int blockId =

blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;

int threadIc =

blockId * (blockDim.x * blockDim.y * blockDim.z)

+ (threadIdx.z *(blockDim.x * blockDim.y))

+ (threadIdx.y *blockDim.x) + threadIdx.x;

内置变量

在编写CUDA代码时，有一些内置变量可以直接使用。内置变量不允许获取变量地址，也不允许被赋值。

CUDA特殊数据类型

Dim3数据类型是CUDA中比较特殊的数据类型，Dim3是基于uint3类型的3维结构体，定义一个dim3类型，默认取值都为1。

CUDA运行库函数

CUDA本身提供了一些可以使用的函数，主要有以下几类

1：设备管理类的函数，例如cudaGetDeviceCount等，这类函数在主机侧被调用，用于查询设备相关的信息。

2：内存管理相关函数，例如cudaMalloc等函数，这类函数允许用户申请设备内存以及在主机内存于设备内存之间进行数据传递。

3：流管理相关函数。用于创建和销毁流。

4：事件管理相关函数，事件的创建和销毁，跟踪等。

5：原子函数等其他函数。

CUDA流与事件

一个典型的CUDA程序，应该是GPU完成一部分工作，CPU完成一部分工作，当CPU把任务交给GPU执行时，有两种策略，第一种策略是等待GPU完成后，CPU继续执行其他任务（同步机制）；另外一种策略是，CPU把任务交给GPU后，不做等待，继续执行需要在CPU侧完成的任务，经过一定时间后，通过查询的方式查询GPU是否已经完成了对应的任务，或者当GPU完成任务后，通知CPU任务已经完成，为了支持第二种策略，CUDA提供给了流和事件机制。

CUDA流

CUDA流可以看作是在GPU上执行任务的一个队列，并且该队列中的操作任务可以按照指定的顺序执行，并且可以创建多个流（多个操作队列），在流和流之间也是可以并行执行。

在CUDA中，流被分为两种类型，隐式流（或者说是匿名流），该流没有名称，无法直接进行控制，在执行上默认与CPU之间是同步模式；显式流（需要主动的声明和创建对应的流），可以对该流进行代码层的直接控制（可以创建同步流，CPU需要等待GPU的执行结果，也可以创建异步流，CPU不需要等待GPU执行结果）。

不同流中的操作可以并行执行，但同一个流中的操作无法并行执行。Hyper-Q技术的出现可以支持多个流实现完全的并行执行，逻辑图如下：

流也可以定义不同的优先级，从而让高优先级的流可以得到更多的执行机会。

在CUDA中，主要通过以下几个接口完成流的创建和使用

1：创建一个流cudaError_t cudaStreamCreate(cudaStream_t* stream);

2：查询该流对应的操作是否已经完成cudaError_t cudaStreamQuery(cudaStream_t stream)

3：cudaError_t cudaStreamSyncronize(cudaStream_t stream);阻塞当前进程执行，直到设备上对应的流中的所有操作都完成。

4：cudaError_t cudaStreamDestroy(cudaStream_t stream); 释放流所占用的资源。

CUDA事件

CUDA中事件有两个作用。

1：可以在流的执行过程中添加标记点，检查执行流是否已经到达对应点，可用于执行过程的等待和测试，作用类似于cudaStreamSyncronize。

2：可以插入到不同的流中，在不同的流中进行操作过程的同步，也即可以通过主机端对设备端执行过程进行操控。

3：可用于执行过程的执行时长统计。

主要有以下几个接口可以使用

1：cudaError_t

cudaEventCreate(cudaEvent_t* event);创建要给事件并得到一个句柄。

2：cudaError_t

cudaEventRecord(cudaEvent_t event, CUstream stream);记录一个事件，当给定了流参数后，当流中所有操作都完成后，该事件被记录。

3：cudaError_t

cudaEventQuery(cudaEvent_t event);查询事件是否已经被记录。

4：cudaError_t cudaEventSyncronize(cudaEvent_t

event);阻塞执行过程，直到对应事件已经发生。cudaError_t cudaEventDestroy(cudaEvent_t event);销毁事件占用的资源。

CUDA代码编译

编写的CUDA程序代码，不能直接使用已有的编译器进行编译，需要使用nvcc工具先进行处理，从编写的代码中分离出设备侧代码，并将设备侧代码编译为二进制格式。

分离出的代码文件可以直接使用标准编译器进行编译，然后于CUDA runtime库链接生成可执行程序。

CUDA代码样例

匿名流样例

一个典型的CUDA程序如下：

1：定义可以从主机侧调用的kernel函数

__global__宏限制addKernel函数在主机侧调用，在GPU侧被执行。传递的参数a和b做为加法运算的“加数”和“被加数”，参数c做为结果。内部计算时，通过threadIdx获取执行addKernel函数的线程，让每个线程计算对应数组中的某一个。

2：通过CUDA运行库函数，获取和设置用于计算的GPU设备

3：调用cuda相关库函数在GPU上分配用于计算的内存空间

4：调用cuda接口函数将主机内存中的数据拷贝到device GPU侧内存中

5：调用kernel函数，在GPU上进行计算

6：调用CUDA接口等待GPU设备侧完成计算，或执行其主机侧代码。该函数会阻塞，直到GPU侧完成计算后才返回。

7：调用CUDA接口将GPU侧执行的结果拷贝会主机内存

8：释放分配的GPU侧内存，释放占用的GPU设备

CUDA