并行编程-CUDA程序设计

一、GPU简介

 GPU（ Graphics Processing Unit，图形处理器 ）是显卡的主要组成部分，有着很强
的浮点运算和并行计算能力，现在GPU已经不仅仅只是应用在图形图像的处理和显示上，
它已经在很多的通用计算领域有了很广泛的应用。
 NVIDIA Tesla（特斯拉）是NVIDIA推出的一个全新的产品系列，主要应用于于广大科
学研究的高性能计算需求。支持向高效利用能源并行计算能力的转化，可以为技术人员
提供专用的计算资源。

1.1 GPU的浮点计算能力和存储带宽

下图表示了GPU的强大的双精度浮点运算能力，从M1060-K80相比于同年CPU产品的差距逐年增加。

 下图表示了GPU的存储带宽，从M1060-K80存储带宽越来越大，相比于同年CPU产品的差距也在逐年增加。

1.2 CPU/GPU的硬件架构比较

 

二、CUDA简介

GPU有着强大的浮点计算能力，为了能够充分的利用CPU和GPU各自的优点，Nvida推出CUDA（Compute Unified Device Architecture，统一计算设备架构）编程模型，使得GPU能够解决复杂的计算问题。
 CUDA体系组成结构

• 驱动
• 开发库（例如cufft、cublas和cudnn等）
• 运行期环境（应用开发接口，基本数据类型定义，内存管理等）

 CUDA特点

• 通用计算不需要映射到图形API
• 提供不同的应用接口（C/C++、Fortran等）
• 单程序、多数据执行模式
• 分为Host Code 和Device Code

 CUDA的应用领域非常的广泛，例如油气勘探、金融分析、雷达仿真、基因分析、地理
信息系统和深度学习等。通过利用GPU的计算能力，通过CUDA并行编程可以使得相关
的应用程序效率有着数倍甚至数百倍的加速。

三、CUDA安装

CUDA需要安装的三个主要文件

• Nvidia driver
• CUDA tookit SDK（lib文件、include文件和bin文件等）
• CUDA samples

 CUDA的安装过程
驱动安装

a) 打开blacklist.conf文件（vim /etc/modprobe.d/blacklist.conf）
b) 在末尾添加blacklist nouveau
c) Ctrl + Alt + F1登录后安装
d) 关闭lightdm（Ctrl + Alt + F1登录后安装）
e) sudo ./cuda_7.0.28_linux_64.run（只安装驱动，安装完之后重启）
CUDA tookit SDK和CUDA samples安装
a) sudo ./cuda_7.0.28_linux_64.run（安装其余部分）

 设置环境变量

sudo gedit ~/.bashrc
在末尾添加 export PATH=$PATH:/usr/local/cuda-12.0/bin
export LD_LIBRARY_PATH=/usr/local/cuda-12.0/lib64:/lib
source ~/.bashrc

 创建cuda.conf文件

sudo touch /etc/ld.so.conf.d/cuda.conf
sudo gedit /etc/ld.so.conf.d/cuda.conf
在里面添加usr/local/cuda-12.0/lib64/
sudo ldconfig

 测试

which nvcc
/usr/local/cuda-12.0/bin/nvcc 

 

CUDA samples

• 位置：~/CUDA-Samples-12-0/NVIDIA_CUDA-12.0_Samples/
• 编译：在上面的目录中直接make，可以编译所有的samples，每个samples文件夹中也可以单独进行编译
• 红色框框出的文件夹对应不同类型的samples，例如设备的性能、有关模拟仿真的简单例子CUDA中的函数库的简单运用等等
• bin文件夹中是samples生成的可执行文件Common文件夹中主要是数据文件、库文件和程序中用到的头文件等等 

四、CUDA编译流程

 五、CUDA程序Makefile的编写

 CPU端源程序

• 假设源程序文件为main.c
• 编译生成main.o文件
• main.o: main.c
• [Tab] g++ -c vector_add.c

GPU端源程序

• 假设源程序文件为kernel.cu
• 编译生成kernel.o文件（需要添加一些头文件等）
• kernel: kernel.cu
• [Tab] nvcc -ccbin g++ -I/usr/local/cuda/samples/commom/inc -c kernel.cu

 链接生成可执行文件

• test: main.o kernel.o
• [Tab] g++ -o test main.o kernel.o -lcuda -lcudart -lstdc++

 删除生成的文件

• clean:
• [Tab] rm –r test *.o

六、CUDA编程整体流程

 CPU处理优势

• 不规则数据结构
• 不可预测存储模式
• 递归算法
• 分支密集型代码
• 单线程程序

 可用存储空间空间较大
 进行硬盘数据的读入与输出
 GPU处理优势

• 规则数据结构
• 密集型数据，数据处理之间相关性较小

 根据CPU和GPU的不同特性通常采用的编程模型为异构编程，CPU处理逻辑性强的处
理和硬盘数据的读入输出操作，而GPU利用其超强的运算能力处理数据密集型的运算。

 Ø CPU串行代码（Host Code）
Ø GPU并行CUDA代码（Device
Code）
Ø Serial Code（Host）
Ø Parallel Kernel Code（device）

七、GPU存储器模型

 寄存器：片上存储器，访存延时小，每个线程私有，首先使用寄存器。
 局部存储器：片外存储器，访存延时大，每个线程私有，寄存器使用完之后会使用局部存储器。
 共享存储器：每个线程块共享，访存延时小。
 全局存储器：片外存储器，访存延时大，整个线程块网格共享，与CPU进行数据交互。

 常量存储器：空间小，支持随机访问，访存速度与命中率有关，只读存储器，通常会存储程序中常用的不需要变化的量。
 纹理存储器：空间较大，支持二维寻址，有缓存机制，访存速度与命中率有关，只读存储器。

八、CUDA线程模型

 Thread：并行的基本单位

• 寄存器和局部存储器为各线程私有

 Thread Block：相互作用的线程组

• 允许线程块内的线程同步
• 线程之间可以进行通信（通过共享存储器）
• 有1维、2维或者3维
• 根据硬件的不同最多可以包含512或者1024个线程

 Grid：一组Thread Block

• 有1维、2维或者3维
• 通过全局存储器进行数据读写

 Kernel：在GPU上执行的核心程序

• 每一个kernel对应一个Grid

 

 线程

• threadIdx.x threadIdx.y threadIdx.z

 线程块

• blockIdx.x blockIdx.y blockIdx.z
• blockDim.x blockDim.y blockDim.z

 线程网格

• gridDim.x gridDim.y gridDim.z

 线程网格中线程索引的计算（以二维举例）

• Idx = blockIdx.x * blockDim.x+threadIdx.x
• Idy = blockIdx.y * blockDim.y +threadIdx.y
• Id = Idy * blockDim.x * gridDim.x + Idx

 线程和线程块在主函数中的定义

• dim3 block(a, b, c)
• dim3 grid(A, B, C)

CUDA编程的基本流程：
① 从文件中读取数据fread(…);
② 在CPU中进行预处理和逻辑操作等
③ 申请Device端的变量空间
cudaMalloc((void **)&d_Data_in, size1);
cudaMalloc((void **)&d_Data_out, size2);
④ 将处理好的数据从Host端拷贝到Device端
cudaMemcpy(d_data_in, h_data_in, size1,
cudaMemcpyHostToDevice);
⑤ 申请线程并调用Kernel函数
⑥ 将处理好的数据从Device端拷贝到Host端
cudaMemcpy(d_data_in, h_data_in, size1,
cudaMemcpyHostToDevice);
⑦ 数据从内存写到文件
fwrite(…); 

九、CUDA编程举例vector add 

 输入变量
 一维向量h_Data_A、一维向量h_Data_B、长度为N
 求解问题描述
计算一维向量h_Data_A和一维向量h_Data_B的和
 输出变量
一维向量h_Data_C，长度为N

 

 在GPU下运行的kernel函数的文件名后缀为.cu

核函数（__global__）
线程索引号的计算
线程私有变量的申请
shared memory的申请等等
设备端函数（__device__），只可以在device端调用，host端不可调用
 在CPU运行的程序文件名可以为.c .cpp .cu
文件的读写
内存的申请
显存（global memory）的申请
CPU和GPU之间数据的交互
GPU端线程的申请
CPU端函数以及kernel函数的调用
内存的释放
显存的释放 

step1:

step2:

step3:

step4:

 step5:

 

step6:

step7:

 对比：

	函数执行 位置：	可以调用的 设备：
__device__ float DeviceFunc ( )	GPU	GPU
__global__ void KernelFunc ( )	GPU	CPU
__host__ float HostFunc ( )	CPU	CPU

十、减少Warps分支 

 什么是Warps
一个block中的每32个线程组成一个warps
这是一种实现的方式，并不是CUDA编程模型中的一部分（在分配线程是可以对线程分配方式进行调整，起到程序优化的作用）
Warps的执行单元是SM
在一个Warps中的threads的执行方式是SIMD
 什么是Warps分支
 在同一个Warps中的线程执行了不同的操作（比如if判断导致的分支）
产生分支之后再一个Warps中的分支会串行进行
过多的Warps分支会导致性能的下降
 Warps分支举例
If(threadIdx.x > 2) { }
在一个Warps中产生了两个不同的操作，thread0, 1, 2执行相同的操作，而剩下
的threads执行另外的操作。

十一、shared memory的运用 

 合理的运用shared memory在进行程序优化时有着非常重要的作用，shared
memory相比于global memory有着非常明显的优势
访存延时非常小，基本可以忽略
可以随机的进行访问
一个block中的threads可以共享，其他的threads不能进行访问，保证了数
据的安全性。
一个block中的threads可以进行同步，在很多需要线程同步的程序中非常有
用。
 shared memory也有这自己的其他的特点
存储的空间不大
只能一个block内的threads共享，在保证数据安全性的同时，在其他的
thread需要用到数据时，需要更多的操作。 

 图像在进行卷积操作是可以运用shared memory进行优化
假如卷积核较大，在进行整个图像卷积时，相同的像素会访问多次。
图像可能较大，可以进行分快处理。
可减少线程访问global memory的次数 

 

 

 注意上图中if括号内表示的是把数据从global memory中读入少shared memory中，在读
的过程中不同的block会读取不同的数据，分三个通道分别读取，在图像的边缘部分，进行
卷积的时候回超出图像，所以需要在其他的部分置为0，也就是上图中的else部分，每个
shared memory中存有计算小块图像卷积的所有信息。 

 十二、CUDA流水线技术

 CUDA stream
声明：cudaStream_t stream0;
创建：cudaStreamCreate(&stream0);
销毁：cudaStreamDestroy(stream0);
 优化策略
Kernel函数和内存拷贝通过并行流水线的方式进行
cudaMemcpyAsync(deviceInput, hostInput, sizef, cudaMemcpyHostToDevice,stream0);
Kernel函数之间通过并行流水线的方式进行
Kernel<<<DimGrid, DimBlock, 0, stream0>>>(deviceInput, deviceOutput, Size);

 十三、利用nvidia的函数库

 利用Nvidia函数库对程序进行加速
Nvidia公司提供了很多有关线性代数，快速傅里叶变换和矩阵求解等函数库，并
且进行了深层次的优化，可以在CUDA编程时直接调用。
矩阵运算的函数库：cublas
快速傅里叶变换的函数库：cufft
有关深度学习的函数库：cudnn
稀疏矩阵库：cuSPARSE
 举例（cublas）
cublas中数据的存储方式为列存储，所以在编程时应该注意，以cublasSegmm
为例。
实现的是C = alpha*A*B + beta*C的功能。
cublasSgemm('t', 't', row_C, col_C, col_A, alpha,A, col_A, B,col_B, beta, C,
row_C);