当前位置: 代码迷 >> CUDA >> 闵大荒之旅(4) - CUDA预热
  详细解决方案

闵大荒之旅(4) - CUDA预热

热度:883   发布时间:2016-04-29 10:42:55.0
闵大荒之旅(四) ---- 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 文件名