CUDA编程基础与实践笔记

 1、GPU硬件简介

   CPU和GPU的显著区别：一块典型的CPU拥有少数几个快速的计算核心，而一块典型的GPU拥有几百到几千个不那么快速的计算核心。CPU中有更多的晶体管用于数据缓存和流程控制，但GPU中有更多的晶体管用于算术逻辑单元。

一块单独的GPU是无法独立地完成所有计算任务的，它必须在CPU的调度下才能完成特定任务。在由CPU和GPU构成的异构计算平台中，通常将起控制作用的CPU称为主机(host)，将起加速作用的GPU称为设备(device)。

FLOPS(floating-point operations per second），表征计算性能的一个重要参数是浮点数运算峰值，即每秒最多能执行的浮点数运算次数。

有效显存带宽：GPU在单位时间内访问设备内存的字节。

2、CUDA中的线程组织

CUDA程序的编译器驱动为: nvcc，在使用nvcc编译器驱动编译.cu文件时，将自动包含必要的CUDA头文件，如<cuda.h>和<cuda_runtime.h>。

编译cuda程序源文件例子：nvcc hello.cu -o hello

GPU只是一个设备，要它工作的话还需要有一个主机给它下达命令，这个主机就是CPU。

一个典型、简单的CUDA程序结构：

int main(void)

{

主机代码

核函数的调用

主机代码

return 0;

}

CUDA核函数的修饰词：__global__

核函数的返回值必须是空类型，即void

主机在调用一个核函数时，必须指明需要在设备中指派多少个线程，否则设备不知道如何工作，比如: hello_from_gpu<<<网络大小, 线程块大小>>>();

cudaDeviceSynchronize()的作用是：同步主机和设备，所以能够促使缓冲区刷新。

核函数中代码的执行方式是“单指令-多线程”，即每一个线程都执行同一串指令。

核函数的内建变量：

gridDim.x：执行配置中变量grid_size的数值

blockDim.x：执行配置中变量block_size的数值

blockIdx.x：指定一个线程在一个网格中的线程块指标，其取值范围是从0到gridDim.x-1

threadIdx.x：指定一个线程在一个线程块中的线程指标，其取值范围是从0到blockDim.x-1

3、简单CUDA程序的基本框架

一个典型的CUDA程序的基本框架：

头文件包含

常量定义（或者宏定义）

C++ 自定义函数和CUDA核函数的声明

int main(void)

{

分配主机与设备内存

初始化主机中的数据

将某些数据从主机复制到设备

调用核函数在设备中进行计算

将某些数据从设备复制到主机

释放主机与设备内存

}

C++自定义函数和CUDA核函数的定义（实现）

在CUDA中，设备内存动态分配由cudaMalloc()函数实现，该函数原型为：

cudaError_t cudaMalloc(void **address, size_t size);

第一个参数address是待分配设备内存的指针。因为内存（地址）本身就是一个指针，所以待分配设备内存的指针就是指针的指针，即双重指针。

函数执行空间标识符：

__global__：核函数，一般由主机调用，在设备中执行。

__device__：设备函数，只能被核函数或其他设备函数调用，在设备中执行

__host__：主机端普通C++函数，在主机中被调用，在主机中执行

4、获得GPU加速的关键

一个CUDA程序能够获得高性能的必要条件有：

（1）数据传输比例较小

（2）核函数的算术强度较高

（3）核函数中定义的线程数目较多

在编写与优化CUDA程序时，一定要想方设法：

（1）减少主机与设备之间的数据传输

（2）提高核函数的算术强度

（3）增大核函数的并行规模

5、CUDA的内存组织

CUDA中设备内存位置：

GPU芯片内：寄存器内存、共享内存

GPU芯片外：全局内存，常量内存，纹理和表面内存，局部内存

全局内存：为核函数提供数据，并在主机与设备及设备与设备之间传递数据。全局内存的生命周期不是由核函数决定的，而是由主机端决定的。

常量内存：是有常量缓存的全局内存，仅有64K。常量内存仅可读、不可写。由于有缓存，常量内存的访问速度比全局内存高。一个使用常量内存的方法是在核函数外面用__constant__定义变量。

纹理内存和表面内存：类似于常量内存 ，但是容量更大。

寄存器：在核函数中定义的不加任何限定符的变量一般来说 就存放于寄存器。寄存器变量仅仅被一个线程可见，也就是说，每一个线程都有一个变量n的副本。寄存器内存在芯片上，是所有内存中访问速度最高的。

局部内存：核函数中定义的不加任何限定符的变量有可能在寄存器中，也有可能在局部内存中。寄存器中放不下的变量，以及索引值不能在编译时就确定的数组，都有可能放在局部内存中。从硬件来看，局部内存只是全局内存的一部分。

共享内存：存在芯片上，具有仅次于寄存器的读写速度，数量也有限。不同于寄存器的是，共享内存对整个线程块可见，其生命周期也与整个线程块一致。

6、共享内存的合理使用

同步机制：__syncthreads()：该函数可保证一个线程块中的所有线程在执行该语句后面的语句之前都完全执行了该语句前面的语句。该函数只针对同一个线程块中的线程，不同线程块中线程的执行次序依然是不确定的。

共享内存：在一个核函数中定义了一个共享内存变量，就相当于在每一个线程块中有了一个该变量的副本。

（1）静态共享内存例子：__shared__ real s_y[128];

（2）动态共享内存例子：

首先，在调用核函数的执行配置中写下第3个参数:

<<<grid_size, block_size, sizeof(real) * block_size>>>

其次，声明动态共享内存：extern __shared__ real s_y[];

共享内存的bank冲突：当同一个线程束内的多个线程试图访问同一个bank中不同层的数据时，就会发生bank冲突。在一个线程束内对同一个bank中的n层数据同时访问将导致n次内存事物，称为发生了n路bank冲突。

7、原子函数的合理使用

在CUDA中，一个线程的原子操作可以在不受其他线程的任何操作的影响下完成对某个数据的一套"读-写-改"操作。

栈内存：存放的时方法以及局部变量，适合少量数据

堆内存：程序产生的对象，比如new 关键字产生的对象实例，适合大量数据。

8、线程束基本函数与协作组

GPU：若干个SM（流多处理器）

warp（线程束）：一个线程块中连续的32个线程

线程块：线程束的整数倍，并总是被分配到一个SM中

线程束内函数：

__ballot_sync()

__all_sync()

__any_sync()

__shfl_sync()

__shfl_up_sync()

__shfl_down_sync()

__shfl_xor_sync()

9、CUDA流

一个CUDA流指的是由主机发出的在一个设备中执行的CUDA操作。

CUDA程序的并行层次主要有两个：一个是核函数内部的并行，一个是核函数外部的并行，CUDA流就属于后者。

为了实现不同的CUDA流之间的并发，主机在向某个CUDA流中发布一系列命令之后必须马上获得程序的控制权，不用等待该CUDA流中的命令在设备中执行完毕。这样，就可以通过主机产生多个相互独立的CUDA流。

核函数执行配置中的流参数：

my_kernel<<<N_grid, N_block, N_shared, stream_id>>>(函数参数);

10、使用统一内存编程

统一内存：是一种逻辑上的概念，它既不是显存，又不是主机的内存，而是一种系统中的任何处理器都可以访问，并能保证一致性的虚拟存储器。这种虚拟存储器是通过CPU和GPU各自内部集成的内存管理单元实现的。在某种程度上，可以将一个系统中某个处理器的内存看成整个统一内存的超级大缓存。

统一内存的好处：

（1）统一内存使CUDA编程更加简单。使用统一内存，将不再需要手动将数据在主机与设备之间传输。

（2）可能会提供比手工移动数据更好的性能。

（3）允许GPU在使用了统一内存的情况下，进行超量分配