CUDA编程简介与深度学习模型部署

1. CUDA编程简介

CUDA是英伟达（NVIDIA）针对外设GPU（Graphic Processing Unit）的一种并行计算架构，全称是Computer Unified Device Architecture， 它包含了CUDA指令集架构（ISA）以及GPU内部的并行计算引擎。 根据厂商可将GPU分为 N 卡（NVIDIA显卡） 和 A 卡（ATI显卡），通常使用的都是N卡。英伟达的显卡系列如下：

1.1 为什么GPU可以加速计算？

在AI兴起的时代，科学计算往往涉及非常高维度的矩阵计算，基于CPU的串行计算难以满足其性能要求，因此基于GPU的并行计算应运而生。显然，并行计算要求同时处理很多很多的数据，这就要求硬件有很多核。因此，GPU可以加速最大的原因是其含有上千个计算核，而CPU的核往往都在个位数。CPU的I/O需要数百上千个周期，串行的计算大部分时间都消耗在I/O上，而GPU则不然。N卡的架构如下图所示，采用了单指令多线程架构（SIMT），上一个线程在进行运算操作时下一个线程就开始I/O操作，类似于指令流水线的形式，当Threads数量足够多时，就可以屏蔽I/O所带来的大量开销。此外，CUDA能够在在应用程序中充分利用CPU和GPU各自的优点，使其应用生态非常友好，受到科研界和产业界的热捧。

1.2 Ubuntu中NVidia驱动、CUDA和CUDNN库安装

NVIDIA显卡驱动安装示例如下：

注：Ctrl+Alt+F1进入文本模式

sudo service lightdm stop
sudo apt-get remove nvidia-*
sudo chmod  a+x NVIDIA-Linux-x86_64-xxx.xx.run
sudo ./NVIDIA-Linux-x86_64-xxx.xx.run -no-x-check -no-nouveau-check -no-opengl-files
sudo service lightdm start

CUDA安装示例如下：

注：安装完建立一个软链接到/usr/locl/cuda

sudo service lightdm stop
sudo sh cuda_x.x.xxx_xxx.xx_linux.run
sudo service lightdm start

CUDNN安装示例如下：

sudo cp cuda/include/cudnn.h /usr/local/cuda/include/ 
sudo cp cuda/lib64/libcudnn* /usr/local/cuda/lib64/ 
sudo chmod a+r /usr/local/cuda/include/cudnn.h 
sudo chmod a+r /usr/local/cuda/lib64/libcudnn*

设置环境：

export PATH=/usr/local/cuda/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH
export CUDA_HOME=$CUDA_HOME:/usr/local/cuda

1.3. CUDA中的Grid、Block与Thread设置

CUDA并行编程的基本思路是把一个很大的任务划分成N个简单重复的操作，创建N个线程分别执行。Thread，block，grid是CUDA编程上的概念，用于组织线程。

thread：一个CUDA的并行程序会被以许多个threads来执行。
block：数个threads会被群组成一个block，同一个block中的threads可以同步，也可以通过shared memory通信。
grid：多个blocks则会再构成grid。

CUDA的软件架构由网格（Grid）、线程块（Block）和线程（Thread）组成，相当于把GPU上的计算单元分为若干（2~3）个网格，每个网格内包含若干（65535）个线程块，每个线程块包含若干（512）个线程，三者的关系如下图：

1.4. CUDA编程说明

（1）存储

在编写CUDA程序时常用的存储结构：

__Host__修饰符为PC主存，即CPU可调度的内存。
__Global__修饰符为Device内存，也就是显存，显卡可以调度的存储空间。
__Shared__修饰符为Shared内存，共享内存，速度与L2 Cache相同级别，且延迟很低，读取周期很短，但要注意Bank Conflict问题。
__Device__修饰符通常用在核函数外，存储在Device内存上，作为全局变量。
__Constant__常量内存。
数据拷贝时，Shared先从 Host 走 PCI-E 通道拷贝到 Device 上，GPU读取指令后先从L1、L2 Cache中寻找数据地址，若没有再从 Device 上寻找，当然 Device 的读取周期漫长。一般的做法是拷贝到 Shared Memory 上，GPU直接从 Shared Memory 取地址调用数据。

（2）调度结构
在调用kernal函数时总体为一个Grid，Grid中含有Block，一个SM在运行时自动分配调用一些Block，每个Block中有大量的Thread。
GPU在运算时以一个Warp为单位，即32个Threads为单位，后面我们可以进行验证调度过程。
Block可以是一维的二维的三维的，Thread也是如此，一般我们选用二维作为调度结构。

（3）性能分析

CUDA程序分析看两个：
1）GFlpos。单浮点数运算峰值性能，刚刚的CUDA-Z已经为我们测算出来了平均性能，我们自己手算一遍峰值性能：840m在默认频率下核心速率为1029MHZ，可以Boosting到1124MHZ（当然不可能一直超频，所以只算默认频率下的计算峰值性能），每颗核心可以同时运行2次乘加，也就是一颗CUDA Core每秒可以运行1029M × 2 = 2058M = 2.058G次乘加，一共有384个CUDA Cores，所以最后每秒可以运行的乘加数为2.058G × 384 = 790.272 GFLOPS。
2）存取带宽。840m的默认显存颗粒频率为900MHZ，等效工作频率为1800MHZ，一次可以取64Bit大小的数据，我们在计算时一般都利用单浮点数进行计算，单浮点数有4个Byte，每个Byte有8个Bit，也就是一个Float型数为32Bit，也就是840m每次操作可以取2个Float型的数，共8Byte大小，故每秒可以取1800 × 8 = 14400MHZ = 14.4GHZ。

2. CUDA编程示例

创建一个CUDA文件，CUDA文件的后缀名是.cu。

首先我们要引入头文件，这里使用Cpp，C，CUDA混编模式，所以需要引入三个头文件：

#include <stdio.h>
#include <iostream>
#include <cuda_runtime.h>
using namespace std;

其中，cuda_runtime.h是cuda函数所在的声明文件。
接下来我们要完成的任务是：将一个矩阵输入到Global内存中，利用GPU全部加1后返回到Host内存进行输出。第一步是需要在CPU中创建一个矩阵，我们一般使用一维动态数组开辟，用二维的方式进行索引。先利用Malloc函数在CPU上开辟一块空间，并全部赋值为1。

int size = 5;
float *a = (float*)malloc(sizeof(float)*size*size);
for (int i = 0; i < size; i++) {
    for (int j = 0; j < size; j++) {
        a[i*size+j] = 1.0f;
    }
}

然后需要在GPU上同样开辟一个相同大小的空间以存放矩阵，这里使用cudaMalloc函数。

float *a_cuda;
cudaMalloc((void**)&a_cuda,sizeof(float)*size*size);

接着，我们将矩阵从CPU上copy到GPU上。

cudaMemcpy(a_cuda,a,sizeof(float)*size*size,cudaMemcpyHostToDevice);

这时的a_cuda指向的是GPU上Device Memory上的一块地址。那GPU要如何才能运行这一块内存中的数据呢？对，没错，就是使用核函数，也叫作Kernel函数。核函数的使用语法如下：

Function<<<griddim,blockdim,extern shared memory,GPU stream>>>(param...);

这里的<<<>>>运算符内是核函数的执行参数，告诉编译器运行时如何启动核函数，用于说明内核函数中的线程数量，以及线程是如何组织的。中间的参数可以控制核函数运行所占用的资源。
griddim表示调用的block块数，blockdim表示调用的thread数，后面两个参数分别表示动态定义共享内存（SM）大小和可使用的SM处理器数。
那说到这里，如何定义kernel呢？kernal函数用__global__修饰符来修饰。下面我们就来定义一个矩阵每个元素都加 1 的kernel函数。在定义核函数之前先要考虑好需要调用多少block和thread，这里时5×5的矩阵，我们可以使用1个block和25个thread排列为5×5thread阵列。核函数定义如下：

__global__ void addone(float *a) {  
    int tix = threadIdx.x;
    int tiy = threadIdx.y;
    int bdx = blockDim.x;
    int bdy = blockDim.y;
    a[tix*bdy+tiy] += 1;
}

其中，threadIdx和blockDim是CUDA的内建变量，指示这当前正在调度的线程号和线程的数量。每个线程都会调用一次，故只需要将a矩阵对应位置的数值+1即可。接着在主函数中调用核函数。

dim3 grid(1, 1, 1), block(5, 5, 1);
addone<<<grid,block>>>(a_cuda);

dim3是一个CUDA内建的变量，是一个三维指示变量，分别对应这x，y，z三维，利用这个变量我们可以控制程序所调用的block和thread总量。由于没有用到动态的shared memory也不去控制调用的SM核心数，所以后面两个参数保持默认值。最后将运行完成的a_cuda拷贝到Host内存中。

cudaMemcpy(a,a_cuda,sizeof(float)*size*size,cudaMemcpyDeviceToHost);

那要如何运行这个程序呢？？？将命令行cd到.cu文件所在目录，利用nvcc编译器进行编译，当然你要知道你的显卡计算能力时多少。整体编译语句风格与GCC类似：

nvcc -gencode=arch=compute_50,code=\"sm_50,compute_50\" -o basic basic.cu

-gencode表示为计算能力xx而生成程序，如果跳出的只有warning而没有error的话说明程序通过编译，可以执行。

./basic

注：(1) 设置thread最好在32的倍数，因为GPU是以warp作为调度单位，设置33这种，实际还是要调用2个warp，实则浪费了31个线程的计算能力。(2) thread并不是开的越多越好，thread少，则程序并行度不够，运算时没有其他的warp进行I/O操作。thread多了，每个SM中寄存器数量有限，thread越多，所能够并行的block就越少，最后还是会导致程序运行缓慢，达不到带宽瓶颈。

3. 神经网络模型部署

在模型部署阶段，一般都会涉及模型加速，这里主要介绍下模型的ONNX与TensorRT转换。

3.1 ONNX模型转换

ONNX是一种神经网络模型保存的中间格式，支持多种格式的模型转为ONNX，也支持使用ONNX导入多种格式的模型，具体见https://github.com/onnx/tutorials。

这里介绍下PyTorch模型转ONNX。在PyTorch下要将模型保存成ONNX格式需要使用torch.onnx.export()函数，使用该函数的时候需要传入下面参数：

• --model: 待保存的model，也就是你在程序中已经训练好或者初始化好的模型
• --input_shape: 指定输入数据的大小，也就是输入数据的形状，是一个包含输入形状元组的列表；
• --name: 模型的名称，即模型的保存路径；
• --verbrose: True或者False,用来指定输出模型时是否将模型的结构打印出来；
• --input_names: 输入数据节点的名称，数据类型为包含字符串的列表；
• --output_names: 输出数据节点的名称，类型与输入数据的节点名称相同；

在成功导出模型后，可以使用ONNX再对模型进行检查：

import onnx

# Load the ONNX model
model = onnx.load("alexnet.onnx")

# Check that the IR is well formed
onnx.checker.check_model(model)

# Print a human readable representation of the graph
onnx.helper.printable_graph(model.graph)

3.2 TensorRT模型转换

TensorRT 是一个深度学习模型线上部署的优化引擎，即 GPU Inference Engine。Tensor 代表张量，即数据流动以张量的方式，如4维张量 [N, C, H, W]。RT表示 runtime。推理性能如图所示，其特性如下：

• 比一些深度学习框架，有更好的优化和及时性
• 支持 custom layer，即用户自定义的神经网络层
• 利用低精度加速推理 inference，如FP16/INT8

这里给出ONNX模型到TensorRT模型的转换与测试代码示例：

def ONNX_build_engine(onnx_file_path):
    '''
    通过加载onnx文件，构建engine
    :param onnx_file_path: onnx文件路径
    :return: engine
    '''
    # 打印日志
    G_LOGGER = trt.Logger(trt.Logger.WARNING)
    with trt.Builder(G_LOGGER) as builder, builder.create_network() as network, trt.OnnxParser(network, G_LOGGER) as parser:
        builder.max_batch_size = 100
        builder.max_workspace_size = 1 << 20
        print('Loading ONNX file from path {}...'.format(onnx_file_path))
        with open(onnx_file_path, 'rb') as model:
            print('Beginning ONNX file parsing')
            parser.parse(model.read())
        print('Completed parsing of ONNX file')

        print('Building an engine from file {}; this may take a while...'.format(onnx_file_path))
        engine = builder.build_cuda_engine(network)
        print("Completed creating Engine")
        
        # 保存计划文件
        # with open(engine_file_path, "wb") as f:
        #     f.write(engine.serialize())
        return engine

def loadONNX2TensorRT(filepath):
    '''
    通过onnx文件，构建TensorRT运行引擎
    :param filepath: onnx文件路径
    '''
    # 计算开始时间
    Start = time()
    engine = self.ONNX_build_engine(filepath)
    # 读取测试集
    datas = DataLoaders()
    test_loader = datas.testDataLoader()
    img, target = next(iter(test_loader))
    img = img.numpy()
    target = target.numpy()

    img = img.ravel()

    context = engine.create_execution_context()
    output = np.empty((100, 10), dtype=np.float32)

    # 分配内存
    d_input = cuda.mem_alloc(1 * img.size * img.dtype.itemsize)
    d_output = cuda.mem_alloc(1 * output.size * output.dtype.itemsize)
    bindings = [int(d_input), int(d_output)]

    # pycuda操作缓冲区
    stream = cuda.Stream()
    # 将输入数据放入device
    cuda.memcpy_htod_async(d_input, img, stream)
    # 执行模型
    context.execute_async(100, bindings, stream.handle, None)
    # 将预测结果从从缓冲区取出
    cuda.memcpy_dtoh_async(output, d_output, stream)
    # 线程同步
    stream.synchronize()

    print("Test Case: " + str(target))
    print("Prediction: " + str(np.argmax(output, axis=1)))
    print("tensorrt time:", time() - Start)

    del context
    del engine

4. NV工具应用

这里主要介绍下nvprof。这是一款用来测试了解并优化CUDA或OpenACC应用程序的性能的分析工具。分析工具使您能够从命令行收集和查看分析数据。

4.1 基本命令

nvprof ./a.out

a.out为编译后的可执行文件，执行结果如下：

注：Profiling result：GPU上运行的时间；API calls：CPU上测量的程序调用API的时间

4.2 相关命令说明

(1) 版本查看

nvprof --version

(2) 总结模式

nvprof aa.py
nvprof ./a.out

其中，a.out是编译生成的可执行文件，aa.py是打算运行的py文件。

3.追踪GPU

nvprof --print-gpu-trace python aa.py
nvprof --print-gpu-trace ./a.out

4.追踪API 

nvprof --print-api-trace python aa.py
nvprof --print-api-trace ./a.out

注：不需要时可以通过–profile-api-trace none关掉这个功能

5. Event/metric总结模式

nvprof --events warps_launched,local_load --metrics ipc  ./a.out

6.Event/metric追踪模式

nvprof --aggregate-mode off --events local_load --print-gpu-trace ./a.out

7.Timeline

nvprof --export-profile timeline.prof ./a.out（python aa.py）
nvprof --metrics achieved_occupancy,executed_ipc -o metrics.prof <app> <app args>
nvprof --kernels <kernel specifier> --analysis-metrics -o analysis.prof <app> <app args>

8.保存为文件

nvprof -o profileOutput ./a.out
nvprof --export-profile timeline.prof ./a.out
nvprof --log-file output.log ./a.out