CUDA并行编程入门：从基础概念到实战技巧

CUDA并行编程入门：从基础概念到实战技巧

引用

CSDN

1.

https://blog.csdn.net/qq_27815483/article/details/140091244

CUDA（Compute Unified Device Architecture）是NVIDIA推出的统一计算设备架构，它使得GPU能够解决复杂的计算问题。本文将从GPU基础、CUDA编程模型、存储器模型、线程模型等多个维度，深入讲解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的硬件架构比较

GPU有着强大的浮点计算能力，为了能够充分的利用CPU和GPU各自的优点，NVIDIA推出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

2.1 驱动安装

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（只安装驱动，安装完之后重启）

2.2 CUDA tookit SDK和CUDA samples安装

sudo ./cuda_7.0.28_linux_64.run（安装其余部分）

2.3 设置环境变量

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

2.4 创建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

2.5 测试

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

2.6 CUDA samples

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

三、CUDA编译流程

四、CUDA程序Makefile的编写

4.1 CPU端源程序

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

4.2 GPU端源程序

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

4.3 链接生成可执行文件

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

4.4 删除生成的文件

• 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编程的基本流程：

1. 从文件中读取数据fread(…);
2. 在CPU中进行预处理和逻辑操作等
3. 申请Device端的变量空间
   cudaMalloc((void **)&d_Data_in, size1);
   cudaMalloc((void **)&d_Data_out, size2);
4. 将处理好的数据从Host端拷贝到Device端
   cudaMemcpy(d_data_in, h_data_in, size1,
   cudaMemcpyHostToDevice);
5. 申请线程并调用Kernel函数
6. 将处理好的数据从Device端拷贝到Host端
   cudaMemcpy(d_data_in, h_data_in, size1,
   cudaMemcpyHostToDevice);
7. 数据从内存写到文件
   fwrite(…);

八、CUDA编程举例vector add

8.1 输入变量

• 一维向量h_Data_A
• 一维向量h_Data_B
• 长度为N

8.2 求解问题描述

计算一维向量h_Data_A和一维向量h_Data_B的和

8.3 输出变量

• 一维向量h_Data_C，长度为N

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

• 核函数（global）
• 线程索引号的计算
• 线程私有变量的申请
• shared memory的申请等等
• 设备端函数（device），只可以在device端调用，host端不可调用

8.5 在CPU运行的程序文件名可以为.c .cpp .cu

• 文件的读写
• 内存的申请
• 显存（global memory）的申请
• CPU和GPU之间数据的交互
• GPU端线程的申请
• CPU端函数以及kernel函数的调用
• 内存的释放
• 显存的释放

8.6 对比

九、减少Warps分支

9.1 什么是Warps

• 一个block中的每32个线程组成一个warps
• 这是一种实现的方式，并不是CUDA编程模型中的一部分（在分配线程是可以对线程分配方式进行调整，起到程序优化的作用）
• Warps的执行单元是SM
• 在一个Warps中的threads的执行方式是SIMD

9.2 什么是Warps分支

• 在同一个Warps中的线程执行了不同的操作（比如if判断导致的分支）
• 产生分支之后再一个Warps中的分支会串行进行
• 过多的Warps分支会导致性能的下降

9.3 Warps分支举例

If(threadIdx.x > 2) { }

在一个Warps中产生了两个不同的操作，thread0, 1, 2执行相同的操作，而剩下的threads执行另外的操作。

十、shared memory的运用

10.1 合理的运用shared memory在进行程序优化时有着非常重要的作用

• 访存延时非常小，基本可以忽略
• 可以随机的进行访问
• 一个block中的threads可以共享，其他的threads不能进行访问，保证了数据的安全性。
• 一个block中的threads可以进行同步，在很多需要线程同步的程序中非常有用。

10.2 shared memory也有这自己的其他的特点

• 存储的空间不大
• 只能一个block内的threads共享，在保证数据安全性的同时，在其他的thread需要用到数据时，需要更多的操作。

10.3 图像在进行卷积操作是可以运用shared memory进行优化

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

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

十一、CUDA流水线技术

11.1 CUDA stream

• 声明：cudaStream_t stream0;
• 创建：cudaStreamCreate(&stream0);
• 销毁：cudaStreamDestroy(stream0);

11.2 优化策略

• Kernel函数和内存拷贝通过并行流水线的方式进行
  cudaMemcpyAsync(deviceInput, hostInput, sizef, cudaMemcpyHostToDevice,stream0);
• Kernel函数之间通过并行流水线的方式进行
  Kernel<<<DimGrid, DimBlock, 0, stream0>>>(deviceInput, deviceOutput, Size);

十二、利用nvidia的函数库

12.1 利用Nvidia函数库对程序进行加速

Nvidia公司提供了很多有关线性代数，快速傅里叶变换和矩阵求解等函数库，并且进行了深层次的优化，可以在CUDA编程时直接调用。

• 矩阵运算的函数库：cublas
• 快速傅里叶变换的函数库：cufft
• 有关深度学习的函数库：cudnn
• 稀疏矩阵库：cuSPARSE

12.2 举例（cublas）

• cublas中数据的存储方式为列存储，所以在编程时应该注意，以cublasSegmm为例。
• 实现的是C = alphaAB + beta*C的功能。
• cublasSgemm('t', 't', row_C, col_C, col_A, alpha,A, col_A, B,col_B, beta, C, row_C);