闵大荒之旅（四） ---- CUDA预热

上一回演示了使用OpenCV中集成的gpu部分进行gpu编程，实现hog+svm算法对行人进行检测，检测效果对比得出gpu运行时间要远比cpu运行时间小，更加具有实时性。但是直接使用OpenCV中的函数，显得过于呆板，灵活性较差，于是，为了将gpu运算牢牢掌握在自己手里，我选择CUDA编程，接下来，我简单的介绍一下gpu、cuda的一部分内容。

 

1.CPU是顺序处理；GPU是并行处理，并行计算将大大缩短时间；

2.CUDA是NVIDIA公司推出的语言风格与C语言很相似的一款通用架构，必须要在NVIDIA公司生产的显卡上使用；

3.CUDA虽然是对GPU并行计算的一个编程框架，但是并不代表这不涉及到CPU：CPU负责进行逻辑性强的事物处理和串行计算，GPU则专注于执行高度线程化的并行处理任务；

4.Jetson TK1板子上并不只是有GPU图形处理器，还有一个CPU（好像是ARM内核的），话说并行计算并不是说光光靠GPU计算就完事了，这只是其中一个部分，最后数据还是要返回到CPU中做最后的处理，只是CPU将比较麻烦、计算量比较大的交给GPU来完成而已！

 

其他的内容我就不详细叙述了，以上几点是我觉得有必要好好了解的。接下来，就开始愉快地开启CUDA编程之旅吧。

先来一段入手程序：test.cu

#include <iostream>
#include <stdio.h>
#include "book.h"
__global__ void kernel(int a, int b, int *c)
{
    *c = a + b;
}
int main()
{
        int c;
        int *dev_c;
        HANDLE_ERROR(cudaMalloc((void**)&dev_c, sizeof(int)));
        kernel<<<1,1>>>(1, 1, dev_c);
        HANDLE_ERROR(cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost));    
        printf("1+1 = %d\n", c);
        cudaFree(dev_c);
        return 0;
}

 乍一看，其实跟C语言几乎没有什么差别，这里只有几个要明白的地方：

1.虽然这个程序看上去跟C语言相似，但是这可不是.C文件，而是.CU文件，可以这么帮助理解，若看到有“__global__”的就可以认定这是个.CU文件，也就是对GPU进行操作的文件；

 

2.头文件：前两个好说，都是大家认识的，而第三个头文件book.h实际上是书《GPU高性能编程CUDA实战》中所附带的头文件，具体下载地址：https://bitbucket.org/mrfright/cuda_by_example/src

 

3.cudaMalloc函数类似于C语言中的开辟内存空间，当我们在GPU内想要得到运行的结果，就要开辟相应的内存空间，否则程序将出错；当然，若是使用到了OpenCV中的GpuMat，那就可以不用写这行代码，因为GpuMat实际上就已经开辟了内存空间（可以理解为内部写好了cudaMalloc吧）

 

4.HANDLE_ERROR（）这个函数是book.h中所写好的，旨在确定（）内程序是否正确执行，若执行错误则抛出相应的错误（其实我觉得不加也行吧）

 

5.__global__ void kernel 实际上就是GPU的核函数，这里我们先介绍一下GPU内部的一些抽象结构：
                          
                     （图片资源来自：http://blog.csdn.net/augusdi/article/details/12439805）

CPU我们称作Host，GPU我们称作Device；

GPU线程结构：

内核以线程网格（Grid）的形式组织，每个线程网格由若干个线程块（block）组成，而每个线程块又由若干个线程（thread）组成。实质上，内核（kernel）是以block为单位执行的，CUDA引入grid只是用来表示一系列可以被并行执行的block的集合，各block的集合。各block是并行执行的，block间无法通信，也没有执行顺序。

所以所以，在__global__声明的这个函数中，实际上就是GPU的内核执行函数，你要做的运算具体内容、具体实现方式都是在这里面完成的

 

6.kernel <<<1,1>>>(1, 1, dev_c);这句话大有名堂啊，实际上也就是调用了5中所说的内核函数，其中：

<<<1,1>>>表示开启了1个block、1个thread线程进行计算，更复杂的情况我将后面阐述；

 

7.cudaMemcpy（源数据， 目标数据， 数据长度， 流向）

cudaMemcpyDeviceToHost表示GPU向CPU传输数据

cudaMemcpyHostToDevice表示CPU想GPU传输数据

 

8.cudaFree表示释放内存，程序员都知道，这是避免溢出的必要手段！

 

运行结果：1 + 1 = 2

表示.CU文件运行成功，成功使用GPU开启线程对1+1进行计算，别看简单，但是打好基础是完成更复杂项目的基础；

 

再来一段代码：

#include "book.h"

#define N 10

__global__ void add(int *a, int *b, int *c)
{
    int tid = blockIdx.x;
    if(tid < N)
    {
        c[tid] = a[tid] + b[tid];
    }
} 


int main()
{
    int a[N], b[N], c[N];
        int *dev_a, *dev_b, *dev_c;
        HANDLE_ERROR(cudaMalloc((void**)&dev_a, N * sizeof(int)));
        HANDLE_ERROR(cudaMalloc((void**)&dev_b, N * sizeof(int)));
        HANDLE_ERROR(cudaMalloc((void**)&dev_c, N * sizeof(int)));
        for(int i = 0; i < N; i++)
        {
        a[i] = i;
                b[i] = i * i;
        }
        HANDLE_ERROR(cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice));
        HANDLE_ERROR(cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice));
        add<<<N, 1>>>( dev_a, dev_b, dev_c);    
        HANDLE_ERROR(cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost));       
 
        //show result
        for (int i = 0; i < N; i++)
        {
        printf("%d + %d = %d\n", a[i], b[i], c[i]);    
    }

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);
    return 0;
}

这段代码相对来说复杂了一点，注意要点如下：

add<<<N, 1>>>相当于开启了10个block每个block一个线程进行运算，实际上也就是数组中的10个元素同时进行运算，那么如何才知道哪个block计算元素的哪个元素呢？

注意到int tid = blockIdx.x;这里我们获取到了block的ID！

因为大家都是执行的相同的指令代码，不同的只是自己的编号，所以，利用这一点，我们可以将数据分开并行运算！是不是很机智！具体内容：

threadIdx，顾名思义获取线程thread的ID索引；如果线程是一维的那么就取threadIdx.x，二维的还可以多取到一个值threadIdx.y，以此类推到三维threadIdx.z。
blockIdx，线程块的ID索引；同样有blockIdx.x，blockIdx.y，blockIdx.z。
blockDim，线程块的维度，同样有blockDim.x，blockDim.y，blockDim.z。
gridDim，线程格的维度，同样有gridDim.x，gridDim.y，gridDim.z。
对于一维的block，线程的threadID=threadIdx.x。
对于大小为（blockDim.x, blockDim.y）的 二维 block，线程threadID=threadIdx.x+threadIdx.y*blockDim.x。
对于大小为（blockDim.x, blockDim.y, blockDim.z）的 三维 block，线程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。

运行结果不用多说了！

ubuntu下编译指令：nvcc 文件名.cu -c -o 文件名