《利用CUDA加速图像处理，提升缺陷检测精度》

课程：利用CUDA加速图像处理，提升缺陷检测精度

CUDA简介¶

CUDA, Compute Unified Device Architecture

CUDA是GPU的一个编程模型
支持C/C++/Python/Fortran/Java/……
全球超过两百五十万开发者，国内超过55万

CUDA C/C++

支持异构编程的扩展方法
简单明了的APIs，能够轻松的管理存储系统

异构计算¶

异构计算

把计算量最大的、最核心的内容，分发给GPU，让GPU帮我们计算
把流程性的代码、逻辑控制的代码，交给CPU来处理

这里涉及到两个术语

【Host】CPU和内存(host memory)
【Device】GPU和显存(device memory)

某年高性能计算大会中

【ISC Top500】137 none GPU；342 GPU
【ISC Top10】3 none GPU；7 GPU

GPU硬件简介¶

基于GPU优化计算时，一个思路就是将串行算法改成并行算法。这种思路的转变，正是因为芯片结构不一样。

GPU计算核心更多，但是Control（控制、流水线的优化）比较少，因此在逻辑处理方面没有CPU强，而计算方面比CPU强
而且GPU计算核心可以共享同一条计算指令，还可以计算同一块Memory

下图为GA100的硬件架构图

绿色的即是计算核心
GPC，Graphics Processing Cluster，图形处理单元集群，GPC由很多SM组成

SM在CUDA中是一个很重要的概念，称之为流多处理器（Streaming Multiprocessor）

L1缓存是内部所有人可以共享的，L0是一个计算单元私有的
在GPU中，多个硬件核心是可以共用一个相同指令集的，32个计算核心称为一个Warp
- 如果你只需调用1个线程，硬件层面也是调用1个Warp（1个线程放到1个核心里计算）
- 32个线程，也是调用1个Warp
- 33个线程，需要调用2个Warp
分发器单元

CUDA安装¶

【适用设备】所有包含NVIDIA GPU的服务器、工作站、个人电脑、嵌入式设备等电子设备

【软件安装】

Windows：安装教程，只需安装一个exe
Linux：安装教程，需要6、7个步骤
Jetson：JetPack SDK | NVIDIA Developer，，利用NVIDIA SDK Manager或者SD Image进行刷机即可

CUDA线程层次¶

三个线程层次¶

CUDA是一个并行计算的编程模型，并行计算途中会同时申请很多个线程。为了方便管理这些线程，我们需要将它们分层次：

	Thread	Thread Block	Thread Grid
	sequential execution unit	a group of threads	a collection of thread blocks
层次	一个线程	多个线程组合成一块，称为block	多个block组成grid
硬件调度	Core级别的调度单位	SM级别的调度单位	GPU（GPC）级别的调度单位
关系	一个线程肯定运行在一个Core中	block中的线程运行在一个SM当中	Grid中的线程运行在一个设备中
	一个Core可能运行多个线程	一个SM可能运行多个block	一个Grid当中的Block可以在多个SM中执行
其他	并行执行	同一个Block中的线程可以协作

硬件调度单位¶

	Threads/Warp	Block(CTA)	Grid
硬件调度	Core级别的调度单位	SM级别的调度单位	GPU（GPC）级别的调度单位
	允许同一个Warp中的thread读取其他thread的值	同一个SM(Streaming Multiprocessor)同一个SM(Shared Memory)	共享同样的Kenrl和Context

【注意】

同CPU的核心一样，当执行一个线程时遇到数据IO，就会产生等待，此时Core就会将此线程挂起，转而去执行另一个线程
同样，一个SM中不仅只有一个block，可能还有很多个block。当一个block遇到数据交换或等待时，SM就会去执行另一个block
同理，一个设备中可以运行多个Grid

【如此】

通常也会将计算核心的忙碌程度，作为衡量程序优化效率的标准
因此，要尽量减少程序中，数据交换、产生分支（如if）、等待的部分

编号¶

【执行设置】dim3 grid(3,2,1), block(5,3,1)

Grid是3x2x1的维度
block是5x3x1的维度

threadIdx.[x,y,z] //是执行当前kernel函数的线程在block中的索引值
blockIdx.[x,y,z]  //是执行当前kernel函数的线程所在block，在grid中的索引值
blockDim.[x,y,z]  //表示一个grid中包含多少个block
gridDim.[x,y,z]   //表示一个block中包含多少个线程

【例子】

int a[32]
blockDim.x==8（1个block有8个线程），y、z维度都为1
如何在CUDA编程中取a[21]元素

需要找到一个线程在全局当中的索引值

//block的y、z维度都为1
int index = threadIdx.x  + blockIdx.x * blockDim.x;
          = 5 + 2 * 8;
          = 21;

示例：向量相加¶

【问题】计算N个向量相加

CPU实现¶

void vecAdd(int n, int* a, int* b, int* c) {
    for(int i{0}; i<n; ++i) c[i] = b[i] + a[i];
}

void main(){
    int size = N * sizeof(int);
    int *a, *b, *c;
    a = (int*)malloc(size);
    b = (int*)malloc(size);
    c = (int*)malloc(size);

    memset(c, 0, size);
    init_rand_f(a, N);
    init_rand_f(b, N);

    vecAdd(N, a, b, c);
}

GPU实现¶

GPU Core并不能直接访问CPU的内存，因此需要将数据从CPU通过PCI传输到GPU当中

一、在GPU中分配内存

二、传输数据

三、每个GPU Core核心从GPU内存中读取、并计算

四、计算完成后，传输回CPU

五、回收GPU内存

《利用CUDA加速图像处理 ，提升缺陷检测精度》