算法优化系列笔记——cuda入门

1 引言

算法在工程化过程中，最躲不开就是算法的优化问题。优化分很多个方向，最简单的实现方式是并行化加速。如：一个向量相加，在cpu中你是串行一个元素一个元素的加减，如果采用并行化加速，你可以一次操作，可以将向量加法完成。当然，并行化的实现方式大概分两类：(1) cpu多线程的方法，如：openmp，(2)异构计算的方法。如gpu加速，fpga加速，NPU加速等。cpu多线程的方法的必要条件就是你得有足够的计算资源，往往在自动驾驶等任务中，cpu往往是稀缺资源。所以异构计算是最常被使用的方法，而异构计算中，没有比N卡的cuda框架更通用的了。为此，我们总结一下cuda的使用方法。

说明：cuda的安装和配置网上资源比较多，暂时先不写了。等以后有时间再补上吧。强调一下gpu驱动版本一定要和cuda版本相对应，不然没法运行

检查下环境

nvcc --version
#输出类似
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2015 NVIDIA Corporation
Built on Tue_Aug_11_14:27:32_CDT_2015
Cuda compilation tools, release 7.5, V7.5.17

我们以一个例子，作为入门课程。而不是一上来就来一堆硬件知识，直接把你干懵逼。

2 cpu版的数组加法

2.1 自动内存管理

cpu_arr_add.cpp

#include <stdio.h>
int main(){
	
	//step1
    int num = 10;
    int a[num],b[num],c[num];
    //step2
    for(size_t i=0; i<num; ++i){
        a[i] = i;
        b[i] = 2*i;
    }
    //step3
	for(size_t i=0; i<num; ++i){
		c[i] = a[i] + b[i];
	}
	
    for(size_t i=0; i<num; ++i){
        printf("%d + %d = %d\n", a[i], b[i], c[i]);
    }
    //step4
    return 0;
}

采用gcc编译一下

gcc cpu_arr_add.cpp -o cpu_arr_add

运行一下

./cpu_arr_add 
0 + 0 = 0
1 + 2 = 3
2 + 4 = 6
3 + 6 = 9
4 + 8 = 12
5 + 10 = 15
6 + 12 = 18
7 + 14 = 21
8 + 16 = 24
9 + 18 = 27

我们来分析一下，整个程序执行的过程。整个执行过程可以分为下面四个步骤：

step4　需要稍微理解一下，a[i],b[i],c[i]在程序执行完后，系统会自动释放。

2.2 手动管理内存

cpu_arr_add_malloc.cpp

#include <stdio.h>
#include <stdlib.h>

int main(){

    int num = 10;
    int *a,*b,*c;
	//step1
    a = (int *) malloc(num * sizeof(int));//c 语言中，手动分配内存
    b = (int *) malloc(num * sizeof(int));
    c = (int *) malloc(num * sizeof(int));
	//step2
    for(size_t i=0; i<num; ++i){
        a[i] = i;
        b[i] = 2*i;
    }
	//step3
	for(size_t i=0; i<num; ++i){
		c[i] = a[i] + b[i];
	}

    for(size_t i=0; i<num; ++i){
        printf("%d + %d = %d\n", a[i], b[i], c[i]);
    }
	//step4
    free(a);//释放内存
    free(b);
    free(c);

    return 0;
}

采用gcc编译一下

gcc cpu_arr_add_malloc.cpp -o cpu_arr_add_malloc

运行一下

./cpu_arr_add_malloc 
0 + 0 = 0
1 + 2 = 3
2 + 4 = 6
3 + 6 = 9
4 + 8 = 12
5 + 10 = 15
6 + 12 = 18
7 + 14 = 21
8 + 16 = 24
9 + 18 = 27

我们再次来分析一下，手动管理内存执行的过程。整个执行过程可以分为下面四个步骤：

3 cuda版数组加法

vector_add.cu

#include <stdio.h>

__global__ void arr_add(int* a,int* b,int* c,int num){
    int i = threadIdx.x;
    if(i<num){
        c[i] = a[i] +b[i];
    }
}

/**
典型的CUDA程序的执行流程如下：

step 1: 分配host内存，并进行数据初始化；
step 2: 分配device内存，并从host将数据拷贝到device上；
step 3: 调用CUDA的核函数在device上完成指定的运算；
step 4: 将device上的运算结果拷贝到host上；
step 5: 释放device和host上分配的内存。

*/
int main(){

    //step 1.1: 分配host内存
    int num = 10;
    int a[num],b[num],c[num];
    int *a_gpu, *b_gpu, *c_gpu;
    //step 1.2: 进行数据初始化
    for(size_t i=0; i<num; ++i){
        a[i] = i;
        b[i] = 2*i;
    }

    //step 2.1: 分配device内存
    cudaMalloc((void **)&a_gpu, num*sizeof(int));
    cudaMalloc((void **)&b_gpu, num*sizeof(int));
    cudaMalloc((void **)&c_gpu, num*sizeof(int));

    //step 2.2: 从host将数据拷贝到device上
    cudaMemcpy(a_gpu,a,num*sizeof(int),cudaMemcpyHostToDevice);
    cudaMemcpy(b_gpu,b,num*sizeof(int),cudaMemcpyHostToDevice);

    //step 3: 调用CUDA的核函数在device上完成指定的运算；
    arr_add<<<1,10>>>(a_gpu,b_gpu,c_gpu,num);

    //step 4: 将device上的运算结果拷贝到host上；
    cudaMemcpy(c,c_gpu,num*sizeof(int),cudaMemcpyDeviceToHost);

    for(size_t i=0; i<num; ++i){
        printf("%d + %d = %d\n", a[i], b[i], c[i]);
    }

    //step 5: 释放device和host上分配的内存
    cudaFree(a_gpu);
    cudaFree(b_gpu);
    cudaFree(c_gpu);
    
    return 0;
}

nvcc vector_add.cu -o vec_add

./vec_add                    
0 + 0 = 0
1 + 2 = 3
2 + 4 = 6
3 + 6 = 9
4 + 8 = 12
5 + 10 = 15
6 + 12 = 18
7 + 14 = 21
8 + 16 = 24
9 + 18 = 27

有没有看到似曾相识的感觉。只是同样内存分配了两次，还有内存拷贝，还有两次内存释放

有没有觉得看不懂，有些内容不认识，如：cudaMalloc,cudaMemcpy,cudaMemcpyHostToDevice
还有__global__，<<<>>>这个三个尖括号是什么鬼

在讲解这个问题前，我们来解析两个概念。我们电脑上现在有两个计算单元，一个CPU和一个GPU。我们来分析一下异同。

内容	CPU	GPU
异构计算中名称	Host	Device
内存分配	malloc	cudaMalloc
内存释放	free	cudaFree
运行方式	串行	并行
标示符	__host__	__global__
编译工具	gcc,g++	nvcc

我们来看下，GPU 数组加法的过程

我们发现主要增加了三项内容，一、给device(GPU)分配内存；二、将数据从host(CPU)拷贝到device(GPU)；三、将计算结果从device拷贝到host。这是因为计算是在GPU上完成的。

我们终于到了，我们神奇的<<<>>>了，这个是干嘛的了？

这就是我们为啥要用gpu加速的计算的原因了，我们前面说了gpu是并行计算，到底怎么实现的了，关键就在这个符号内传入的参数。arr_add<<<1,10>>>(a_gpu,b_gpu,c_gpu,num);在我们的程序中，我相当于用了1x10的线程并行计算。而这两个参数到底是怎么设置的，我们下个笔记详细说明。

4 总结分析

我们来分析下cpu和gpu运算之间的区别

//cpu中
	//step3
	for(size_t i=0; i<num; ++i){
		c[i] = a[i] + b[i];
	}

//gpu中
__global__ void arr_add(int* a,int* b,int* c,int num){
    int i = threadIdx.x;
    if(i<num){
        c[i] = a[i] +b[i];
    }
}
    //step 3: 调用CUDA的核函数在device上完成指定的运算；
    arr_add<<<1,10>>>(a_gpu,b_gpu,c_gpu,num);

在cpu中我们用for循环，执行完一圈然后回去再执行下一圈
在gpu中我调用了__global__ void arr_add函数，发现for循环消失了。而是同时开了10个线程同时进行计算。

5 拔

我们主要解释下cuda计算的主要思想和第一个能完全运行起来的程序。在实践中学习，才是进步最快的，在下一篇笔记中，我们将解释下各个参数设置的技巧，以及我们如何在程序中使用这些加速算法。