传统的中央处理器（CPU，Central Processing Unit) 内部结构异常复杂，主要是因为其需要很强的通用性来处理各种不同的数据类型，同时又要逻辑判断又会引入大量的分支跳转和中断的处理。 为了提高计算能力，CPU通常会采取提高时钟频率或增加处理器核数量的策略。为了进一步获得更高效的计算，图形处理器（GPU, Graphics Processing Unit）应运而生。 GPU可以在无需中断的纯净环境下处理类型高度统一的、相互无依赖的大规模数据。

GPU的高效在于可以高度并行处理。 以两个向量相加为例，CPU可能采取循环处理，每个循环对一个分量做加法。GPU则可以开多个线程，每个线程同时对一个分量做加法。CPU加法的速度一般快于GPU，但因为GPU可以同时开大量线程并行跑，因此更加高效。为了降低GPU程序的开发难度，NVIDIA推出了 CUDA（Compute Unified Device Architecture，统一计算设备架构）这一编程模型。

一. 环境配置

在保证NVIDIA显卡驱动成功安装的条件下，从下面链接下载并安装对应版本的CUDA Toolkit（注意：最好已经安装好VS), 通过在命令窗中执行 nvcc -V初步判断是否安装成功：

nvcc -V

安装成功后(默认安装)系统会增加如下环境变量

CUDA_PATH：  C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0
CUDA_PATH_V8_0：  C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0
NUMBER_OF_PROCESSORS： 4
NVTOOLSEXT_PATH：  C:\Program Files\NVIDIA Corporation\NvToolsExt\

CUDA Toolkit安装成功后会自动和系统的编译器进行绑定。 “新建项目”下增加了 “NVIDIA”选项。

之后发现项目里多了一个 “kernel.cu”的文件，该文件内容是一个经典的 矢量相加 的GPU程序。 CUDA C编程和我们熟悉的标准C在很大程度上是没有区别的。 同时，这段程序直接运行在 主机上。

接下来，我们看看如何使用GPU来执行代码。如下：

__global__  void mkernel(void){}

int main()
{
    mkernel <<<1,1>>>();
    std::cout<<"Hello, World!"<<std::endl;
    system("pause");
    return 0;
}

这里主要有一个空的函数mkernel()， 并带有修饰符 __global__ 对空函数的调用， 并带有修饰符 <<<1,1>>>_global_ 为CUDA C为标准C增加的修饰符，表示该函数将会交给编译设备代码的编译器(NVCC)并最终在设备上运行。 而 main函数则依旧交给系统编译器(vs)。其实，CUDA就是通过直接提供API接口或者在语言层面集成一些新的东西来实现在主机代码中调用设备代码。

二. CUDA核函数运行参数

2.1 核函数运行参数

当我们使用 gloabl 声明核函数后

__global__ void kernel(param list){  }

在主机端(Host)调用时采用如下的形式：

kernel<<<Dg, Db, Ns, S>>>(param list);

• Dg：int型或者dim3类型(x,y,z)。 用于定义一个grid中的block是如何组织的。 int型则直接表示为1维组织结构。
• Db int型或者dim3类型(x,y,z)。 用于定义一个block中的thread是如何组织的。 int型则直接表示为1维组织结构。
• Ns size_t类型，可缺省，默认为0。 用于设置每个block除了静态分配的共享内存外，最多能动态分配的共享内存大小，单位为byte。 0表示不需要动态分配。
• S：cudaStream_t类型，可缺省，默认为0。 表示该核函数位于哪个流。

2.2 线程结构

关于CUDA的线程结构，有着三个重要的概念： Grid, Block, Thread

• GPU工作时的最小单位是 thread。
• 多个 thread 可以组成一个 block，但每一个 block 所能包含的 thread 数目是有限的。因为一个block的所有线程最好应当位于同一个处理器核心上，同时共享同一块内存。 于是一个 block中的所有thread可以快速进行同步的动作而不用担心数据通信壁垒。
• 执行相同程序的多个 block，可以组成 grid。 不同 block 中的 thread 无法存取同一块共享的内存，无法直接互通或进行同步。因此，不同 block 中的 thread 能合。不过，利用这个模式，可以让程序不用担心显示芯片实际上能同时执行的 thread 数目限制。例如，一个具有很少量执行单元的显示芯片，可能会把各个 block 中的 thread 顺序执行，而非同时执行。不同的 grid 则可以执行不同的程序(即 kernel)。

下图是一个结构关系图：

此外，Block, Thread的组织结构可以是可以是一维，二维或者三维。以上图为例，Block, Thread的结构分别为二维和三维。

CUDA中每一个线程都有一个唯一标识ThreadIdx，这个ID随着组织结构形式的变化而变化。 (注意：ID的计算，同计算行优先排列的矩阵元素ID思路一样。)

// Block是一维的，Thread也是一维的
__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = blockIdx.x *blockDim.x + threadIdx.x;  
    c[i] = a[i] + b[i];
}

// Block是一维的，Thread是二维的
__global__ void addKernel(int *c, int *a, int *b)
{
    int i = blockIdx.x * blockDim.x * blockDim.y + \
        threadIdx.y * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

// Block是二维的，Thread是三维的
__global__ void addKernel(int *c, int *a, int *b)
{
    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int i = blockId * (blockDim.x * blockDim.y * blockDim.z)  
        + (threadIdx.z * (blockDim.x * blockDim.y))  
        + (threadIdx.y * blockDim.x) + threadIdx.x; 
    c[i] = a[i] + b[i];
}

当然也可以通过下面的代码来直接查询自己GPU的具体指标：

cudaError_t cudaStatus;

// 初获取设备数量
int num = 0;
cudaStatus = cudaGetDeviceCount(&num);
std::cout << "Number of GPU: " << num << std::endl;

// 获取GPU设备属性
cudaDeviceProp prop;
if (num > 0)
{
    cudaGetDeviceProperties(&prop, 0);
    // 打印设备名称
    std::cout << "Device: " <<prop.name << std::endl;
}

其中 cudaDeviceProp是一个定义在driver_types.h中的结构体，大家可以自行查看其定义。

2.3 内存结构

如下图所示,每个 thread 都有自己的一份 register 和 local memory 的空间。同一个 block 中的每个 thread 则有共享的一份 share memory。此外，所有的 thread(包括不同 block 的 thread)都共享一份 global memory、constant memory、和 texture memory。不同的 grid 则有各自的 global memory、constant memory 和 texture memory。

三. 线程协作

3.1 并行程序块的分解

首先回顾我们之前实现的矢量相加程序：

// 核函数：每个线程负责一个分量的加法
__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x; // 获取线程ID
    c[i] = a[i] + b[i];
}

// 运行核函数，运行设置为1个block，每个block中size个线程
addKernel << <1, size >> >(dev_c, dev_a, dev_b);

我们知道一个Block中的可开辟的线程数量是有限的(不超过1024)。如果矢量特别长，上面的操作是会出现问题的。于是我们可以采用多个线程块(Block)来解决线程不足的问题。 假如我们设定每个线程块包含128个线程，则需要的线程块的数量为 size / 128。 为了避免不能整除带来的问题，我们可以稍微多开一点 (size + 127) / 128，但需要增加判断条件来避免越界。

// 核函数：每个线程负责一个分量的加法
__global__ void addKernel(int *c, const int *a, const int *b, const int size)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x; // 获取线程ID
    if (i < size)
        c[i] = a[i] + b[i];
}

// 运行核函数，运行设置为多个block，每个block中128个线程
addKernel <<<(size + 127) / 128, 128 >>>(dev_c, dev_a, dev_b, size);

通过前面小节，我们同时也知道一个Grid中可开辟的Block数量也是有限的。

如果数据量大于 Block_num * Thread_num，那么我们就无法为每个分量单独分配一个线程了。 不过，一个简单的解决办法就是在核函数中增加循环。

// 核函数：每个线程负责多个分量的加法
__global__ void addKernel(int *c, const int *a, const int *b, const int size)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x; 
    while (i < size)
    {
        c[i] = a[i] + b[i];
        // 偏移分量等于一个Grid中包含的线程数量
        i += blockDim.x * gridDim.x;
    }
}

// 运行核函数，运行设置为1个Grid包含128个block，每个block包含128个线程
// 其中已经假设 size > 128*128
addKernel <<<128, 128 >>>(dev_c, dev_a, dev_b, size);

3.2 共享内存与同步

上面提到线程块的分解似乎是为了增加可用的线程数量，但这种说法并不靠谱，因为这完全可以由CUDA在幕后全权控制。 其实，分解线程块的重要原因是因为内存。

在“4.3 内存结构”中我们已经知道，同一个Block中的线程可以访问一块共享内存。由于共享内存缓冲区驻留在物理GPU上，而不是GPU之外的系统内存上，因此访问共享内存的延迟要远远低于访问普通缓冲区的延迟。

不同Block之间存在隔离，如果我们需要不同线程之间进行通信，那么还需要考虑线程同步的问题。比如线程A将某个数值写入内存，然后线程B会对该数值进行一些操作，那么显然必须等A完成之后B才可以操作，如果没有同步，程序将会因进入“竞态条件”而产生意想不到的错误。

接下来我们通过一个矢量点积的例子来说明上述问题。

矢量点积的定义如下：

由上面的定义来看，点积的实现可以分为两步：

• （1）计算每个分量的乘积，并暂存该结果；
• （2）将所有临时结果加和。

// 定义我们的矢量长度
const int N = 64 * 256; 

// 定义每个Block中包含的Thread数量 
const int threadsPerBlock = 256;  

// 定义每个Grid中包含的Block数量, 这里32 < 64， 是为了模拟线程数量不足的情况
const int blocksPerGrid = 32;

__global__ void dot( float *a, float *b, float *c ) 
{  
    // 声明共享内存用于存储临时乘积结果，内存大小为1个Block中的线程数量
    __shared__ float cache[threadsPerBlock];  

    // 线程索引
    int tid = threadIdx.x + blockIdx.x * blockDim.x;  

    // 一个Block中的线程索引 
    int cacheIndex = threadIdx.x;  

    // 计算分量乘积，同时处理线程不足的问题
    float   temp = 0;  
    while (tid < N) {  
        temp += a[tid] * b[tid];  
        tid += blockDim.x * gridDim.x;  
    }  

    // 存储临时乘积结果
    cache[cacheIndex] = temp; 
}

执行完上面的部分，我们剩下的就是把cache中的值相加求和。 但是，我们必须要保证所有乘积都已经计算完成，才能去计算求和。 命令如下：

// 对线程块中的所有线程进行同步
// 线程块中的所有线程都执行完前面的代码后才会继续往后执行
__syncthreads();

求和最直接的方式就是循环累加，此时复杂度与数组长度成正比。不过我们可以再用一种更加高效的方法，其复杂度与数组长度的log成正比：将值两两合并，于是数据量减小一半，再重复两两合并直至全部计算完成。代码如下：

// 合并算法要求长度为2的指数倍
int i = blockDim.x/2;  
while (i != 0) 
{  
    if (cacheIndex < i)  
        cache[cacheIndex] += cache[cacheIndex + i];  
    __syncthreads();  
    i /= 2;  
} 

// 最后将一个Block的求和结果进行保存
if (cacheIndex == 0)  
        c[blockIdx.x] = cache[0]; 

四、 函数与变量类型限定符

4.1 函数类型限定符

函数类型限定符用来标识函数运行在主机(CPU)还是设备(GPU)上，函数由主机还是设备调用。

__global__

• __global__修饰的函数为 核函数。
• 运行在设备上
• 可以由主机调用
• 可以由计算能力大于3.2的设备调用
• 必须有void返回类型
• 调用时必须制定运行参数(<<< >>>)
• 该函数的调用时异步的，即可以不必等候该函数全部完成，便可以在CPU上继续工作

__device__

• 运行在设备上
• 只能由设备调用
• 编译器会内联所有认为合适的__device__修饰的函数

__host__

• 运行在主机上
• 只能由主机调用
• 效果等同于函数不加任何限定符
• 不能与__global__共同使用， 但可以和__device__联合使用

__host__ __device__能够用来修饰类的成员函数，__global__不能（类的静态函数也不行）
__global__只能修饰普通的全局函数。

__noinline__

• 声明不允许内联

__forceinline__

• 强制编译器内联该函数

4.2 变量类型限定符

变量类型限定符用来标识变量在设备上的内存位置。

__device__ (单独使用时)

• 位于 global memory space
• 生命周期为整个应用期间(即与application同生死)
• 可以被grid内的所有threads读取
• 可以在主机中由以下函数读取
  • cudaGetSymbolAddress()
  • cudaGetSymbolSize()
  • cudaMemcpyToSymbol()
  • cudaMemcpyFromSymbol()

__constant__

• 可以和 __device__ 联合使用
• 位于 constant memory space
• 生命周期为整个应用期间
• 可以被grid内的所有threads读取
• 可以在主机中由以下函数读取
  • cudaGetSymbolAddress()
  • cudaGetSymbolSize()
  • cudaMemcpyToSymbol()
  • cudaMemcpyFromSymbol()

__shared__

• 可以和 device 联合使用
• 位于一个Block的shared memory space
• 生命周期为整个Block
• 只能被同一block内的threads读写

__managed__

• 可以和 __device__ 联合使用
• 可以被主机和设备引用，主机或者设备函数可以获取其地址或者读写其值
• 生命周期为整个应用期间

__restrict__

• 该关键字用来对指针进行限制性说明，目的是为了减少指针别名带来的问题。

五、 常量内存与事件

GPU通常包含大量的数学计算单元，因此性能瓶颈往往不在于芯片的数学计算吞吐量，而在于芯片的内存带宽，即有时候输入数据的速率甚至不能维持满负荷的运算。 于是我们需要一些手段来减少内存通信量。 目前的GPU均提供了64KB的常量内存，并且对常量内存采取了不同于全局内存的处理方式。 在某些场景下，使用常量内存来替换全局内存可以有效地提高通信效率。

常量内存

常量内存具有以下特点：

• 需要由 __constant__ 限定符来声明
• 只读
• 硬件上并没有特殊的常量内存块，常量内存只是只是全局内存的一种虚拟地址形式
• 目前的GPU常量内存大小都只有64K，主要是因为常量内存采用了更快速的16位地址寻址(2^16)
• 对于数据不太集中或者重用率不高的内存访问，尽量不要使用常量内存，否则甚至会慢于使用全局内存
• 常量内存无需cudaMalloc()来开辟，而是在声明时直接提交一个固定大小，比如 __constant__ float mdata[1000]
• 常量内存的赋值不能再用cudaMemcpy()，而是使用cudaMemcpyToSymbol()

常量内存带来性能提升的原因主要有两个：

• 对常量内存的单次读操作可以广播到其他的“邻近(nearby)”线程，这将节约15次读取操作
• 常量内存的数据将缓存起来，因此对于相同地址的连续操作将不会产生额外的内存通信量。

对于原因1，涉及到 线程束(Warp)的概念。

在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“步调一致(Lockstep)”的形式执行。 即线程束中的每个线程都将在不同数据上执行相同的指令。

当处理常量内存时，NVIDIA硬件将把单次内存读取操作广播到每个半线程束(Half-Warp)。在半线程束中包含16个线程，即线程束中线程数量的一半。如果在半线程束中的每个线程从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求并在随后将数据广播到每个线程。如果从常量内存中读取大量数据，那么这种方式产生的内存流量只是使用全局内存时的1/16。

对于原因2，涉及到缓存的管理

由于常量内存的内容是不发生变化的，因此硬件将主动把这个常量数据缓存在GPU上。在第一次从常量内存的某个地址上读取后，当其他半线程束请求同一个地址时，那么将命中缓存，这同样减少了额外的内存流量。

另一方面, 常量内存的使用也可能会对性能产生负面的影响。半线程束广播功能实际上是一把双刃剑。虽然当所有16个线程都读取相同地址时，这个功能可以极大提升性能，但当所有16个线程分别读取不同的地址时，它实际上会降低性能。因为这16次不同的读取操作会被串行化，从而需要16倍的时间来发出请求。但如果从全局内存中读取，那么这些请求会同时发出。

六、 原子操作

原子操作 是指对全局和共享内存中的32位或者64位数据进行 “读取-修改-覆写”这一操作。

原子操作可以看作是一种最小单位的执行过程。 在其执行过程中，不允许其他并行线程对该变量进行读取和写入的操作。 如果发生竞争，则其他线程必须等待。

下面先给出原子操作函数的列表，后续会给出一个应用例子。

原子操作函数列表

6.1 atomicAdd()

int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address, unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address, unsigned long long int val);
float atomicAdd(float* address, float val);
double atomicAdd(double* address, double val);

读取位于全局或共享存储器中地址address处的32位或64位字old，计算(old + val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

注意：

• 32位浮点数的操作只适用于计算能力大于2.0的GPU
• 64位浮点数的操作只适用于计算能力大于6.0的GPU

但可以通过以下操作在计算能力不足的GPU上实现浮点数原子操作：

#if __CUDA_ARCH__ < 600 
__device__ double atomicAdd(double* address, double val) 
{ 
    unsigned long long int* address_as_ull = (unsigned long long int*)address; unsigned long long int old = *address_as_ull, assumed; 
    do { 
        assumed = old; 
        old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed))); 
        // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN) 
        } 
        while (assumed != old); 
        return __longlong_as_double(old); 
} 
#endif

6.2 atomicSub()

int atomicSub(int* address, int val); 
unsigned int atomicSub(unsigned int* address, unsigned int val);

读取位于全局或共享存储器中地址address处的32位字old，计算(old - val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

6.3 atomicExch()

int atomicExch(int* address, int val); 
unsigned int atomicExch(unsigned int* address, unsigned int val); 
unsigned long long int atomicExch(unsigned long long int* address, unsigned long long int val); 
float atomicExch(float* address, float val);

读取位于全局或共享存储器中地址address处的32位或64位字old，并将val 存储在存储器的同一地址中。这两项操作在一次原子事务中执行。该函数将返回old。

6.4 atomicMin()

int atomicMin(int* address, int val); 
unsigned int atomicMin(unsigned int* address, unsigned int val); 
unsigned long long int atomicMin(unsigned long long int* address, unsigned long long int val);

读取位于全局或共享存储器中地址address处的32位字或64位字old，计算old 和val 的最小值，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old>

注意：

64位的操作只适用于计算能力大于3.5的GPU

6.5 atomicMax()

同atomicMin()。

6.6 atomicInc()

unsigned int atomicInc(unsigned int* address, unsigned int val);

读取位于全局或共享存储器中地址address处的32位字old，计算 ((old >= val) ? 0 : (old+1))，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

6.7 atomicDec()

unsigned int atomicDec(unsigned int* address, unsigned int val);

读取位于全局或共享存储器中地址address处的32位字old，计算 (((old == 0) | (old > val)) ? val : (old-1))，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

6.8 atomicCAS()

int atomicCAS(int* address, int compare, int val); 
unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val); 
unsigned long long int atomicCAS(unsigned long long int* address, unsigned long long int compare, unsigned long long int val);

读取位于全局或共享存储器中地址address处的32位或64位字old，计算 (old == compare ? val : old)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old（比较并交换）。

6.9 atomicAnd()

int atomicAnd(int* address, int val); 
unsigned int atomicAnd(unsigned int* address, unsigned int val); 
unsigned long long int atomicAnd(unsigned long long int* address, unsigned long long int val);

读取位于全局或共享存储器中地址address处的32位字或64位字old，计算 (old & val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

注意：

64位的操作只适用于计算能力大于3.5的GPU

6.10 atomicOr()

int atomicOr(int* address, int val); 
unsigned int atomicOr(unsigned int* address, unsigned int val); 
unsigned long long int atomicOr(unsigned long long int* address, unsigned long long int val);

读取位于全局或共享存储器中地址address处的32位字或64位字old，计算 (old | val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

注意：

64位的操作只适用于计算能力大于3.5的GPU

6.11 atomicXor()

int atomicXor(int* address, int val); 
unsigned int atomicXor(unsigned int* address, unsigned int val); 
unsigned long long int atomicXor(unsigned long long int* address, unsigned long long int val);

读取位于全局或共享存储器中地址address处的32位字或64位字old，计算 (old ^ val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

注意：

64位的操作只适用于计算能力大于3.5的GPU

七. 流并行

我们前面学习的CUDA并行程序设计，基本上都是在一批数据上利用大量线程实现并行。 除此之外， NVIDIA系列GPU还支持另外一种类型的并行性 —— 流。

GPU中的流并行类似于CPU上的任务并行，即每个流都可以看作是一个独立的任务，每个流中的代码操作顺序执行。

下面从流并行的基础到使用来说明。

7.1 页锁定内存

流并行的使用需要有硬件支持：即必须是支持设备重叠功能的GPU。

通过下面的代码查询设备是否支持设备重叠功能：

cudaDeviceProp mprop;
cudaGetDeviceProperties(&mprop,0);
if (!mprop.deviceOverlap)
{
    cout << "Device not support overlaps, so stream is invalid!" << endl;
}

只有支持设备重叠，GPU在执行一个核函数的同时，才可以同时在设备与主机之间执行复制操作。 当然，这种复制操作需要在一种特殊的内存上才可以进行 —— 页锁定内存。

页锁定内存： 需要由cudaHostAlloc()分配，又称为固定内存（Pinned Memory）或者不可分页内存。 操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中，因为这块内存将不会被破坏或者重新定位。 由于gpu知道内存的物理地址，因此可以通过“直接内存访问（Direct Memory Access，DMA）” 直接在gpu和主机之间复制数据。
可分页内存： malloc()分配的内存是标准的、可分页的（Pagable）主机内存。 可分页内存面临着重定位的问题，因此使用可分页内存进行复制时，复制可能执行两次操作：从可分页内存复制到一块“临时”页锁定内存，然后从页锁定内存复制到GPU。
虽然在页锁定内存上执行复制操作效率比较高，但消耗物理内存更多。因此，通常对cudaMemcpy()调用的源内存或者目标内存才使用，而且使用完毕立即释放。

7.2 流并行机制

流并行是指我们可以创建多个流来执行多个任务， 但每个流都是一个需要按照顺序执行的操作队列。 那么我们如何实现程序加速？ 其核心就在于，在页锁定内存上的数据复制是独立于核函数执行的，即我们可以在执行核函数的同时进行数据复制。

这里的复制需要使用cudaMemcpyAsync()，一个以异步执行的函数。调用cudaMemcpyAsync()时，只是放置一个请求，表示在流中执行一次内存复制操作。当函数返回时，我们无法确保复制操作已经结束。我们能够得到的保证是，复制操作肯定会当下一个被放入流中的操作之前执行。(相比之下，cudaMemcpy()是一个同步执行函数。当函数返回时，复制操作已完成。)

以计算 a + b = c为例，假如我们创建了两个流，每个流都是按顺序执行：

复制a(主机到GPU) -> 复制b(主机到GPU) -> 核函数计算 -> 复制c(GPU到主机)

如上图，复制操作和核函数执行是分开的，但由于每个流内部需要按顺序执行，因此复制c的操作需要等待核函数执行完毕。 于是，整个程序执行的时间线如下图：(箭头表示需要等待)

从上面的时间线我们可以启发式的思考下：如何调整每个流当中的操作顺序来获得最大的收益？ 提高重叠率。

如下图所示，假如复制一份数据的时间和执行一次核函数的时间差不多，那么我们可以采用交叉执行的策略：

由于流0的a和b已经准备完成，因此当复制流1的b时，可以同步执行流0的核函数。 这样整个时间线，相较于之前的操作很明显少掉了两块操作。

7.3 流并行示例

与流相关的常用函数如下

// 创建与销毁
cudaStream_t stream//定义流 
cudaStreamCreate(cudaStream_t * s)//创建流 
cudaStreamDestroy(cudaStream_t s)//销毁流 

//同步 
cudaStreamSynchronize()//同步单个流：等待该流上的命令都完成 
cudaDeviceSynchronize()//同步所有流：等待整个设备上流都完成 
cudaStreamWaitEvent()//等待某个事件结束后执行该流上的命令 
cudaStreamQuery()//查询一个流任务是否完成 

//回调 
cudaStreamAddCallback()//在任何点插入回调函数 

//优先级 
cudaStreamCreateWithPriority() 
cudaDeviceGetStreamPriorityRange()

附录

• NVIDIA 型号对应的计算能力(Compute Capabilities)表

Quadro Desktop Products

Quadro Mobile Products