CUDA系列之：how to start...

练习CUDA基础编程，从1开始。

https://zhuanlan.zhihu.com/p/27916653

此方法简单来说是利用了目前计算机为多核式结构（四核），因为在一般的串行编程中，一个程序的指令在单一的 CPU 上按照先后顺序依次执行，而并行编程则将一个程序分成独立的若干部分在一个或多个 CPU 上进行同步执行以取得更高的运算效率和性能。在我自己的项目中我将图像检测的算法写入一个函数中（算法实现较为简单），通过上述方法，同时对3张传入的待检测图进行识别（摄像机拍摄到的蚕茧图片），部分代码如下

    Mat img1 = imread("new.bmp");//读入第一张图片
    Mat img2 = imread("new1.bmp");//读入第二张图片
    Mat img3 = imread("new2.bmp");//读入第三张图片
    Mat imgResult1, imgResult2, imgResult3;//创建识别结果
    //多线程运算
    auto f1 = async(Pod, img1);//Pod为检测算法函数
    auto f2 = async(Pod, img2);
    auto f3 = async(Pod, img3);

    imgResult1 = f1.get();
    imgResult2 = f2.get();
    imgResult3 = f3.get();

但由于CPU加速效果的有限，在处理一些大数据的时候，这种并行的方式并不能从本质上解决程度处理速度慢的问题，故在5月份有幸接触到了GPU并行加速这一概念，开始学习Nvidia公司推出的基于CUDA的并行程序设计，到目前本人已经完成了一些基于OPENCV的图像处理算法关于CUDA的改编，在加速图像处理程序上有着不错的效果，故打算将自己的一些关于CUDA的学习以及实践进行记录，如有不当之处，希望各位大佬们批评指正~

CUDA编程基础

了解电脑的同学都知道现在每台计算机在其硬件构成上都包括最重要的两部分——主机和显卡，而在CUDA的架构中引入了主机端（host）和设备（device）的概念，主机端即为CPU，而设备端则为GPU，所谓的CUDA编程即为利用GPU在大规模并行计算上的执行效率高的优势（一些对显卡要求高的大型游戏本质上来讲就是要求显卡的计算速度快），利用CPU，将一些逻辑复杂性较小的计算部分分配给GPU让其协助完成计算任务。

一个完整的CUDA程序由主机代码和设备代码两部分组成。主机端代码部分在CPU上执行，是普通的C代码；设备端代码部分在GPU上执行，此代码部分在内核kernel上遍写（.cu文件）。关于CUDA中C语言的编写规范以及一些功能网上已有许多健全的资料，在这里不在多作解释，简单利用一个CUDA程序进行作为例程进行实践：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdlib.h>
#include <iostream>
using namespace std;
//将数组A与B中的元素相加存入数组C中
__global__ void arrayAdd(int *A, int *B, int *C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    //简单的示例，A、B、C均为长度为4的整型数组
    //定义主机端数组A、B、C
    int A[4] = { 1, 2, 3, 4 };
    int B[4] = { 2, 3, 4, 5 };
    int C[4];
    //定义设备端数组d_a、d_b、d_c
    int *d_a;
    int *d_b;
    int *d_c;
    //定义他们在设备端的空间大小
    int size = 4 * sizeof(int);
    //在GPU中为他们开辟分配对应的显存空间
    cudaMalloc((void**)&d_a, size);
    cudaMalloc((void**)&d_b, size);
    cudaMalloc((void**)&d_c, size);
    //利用cudaMemcpy函数在CPU端A,B的值复制到对应的GPU内存中
    cudaMemcpy(d_a, A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, B, size, cudaMemcpyHostToDevice);
    //调用编写好的内核程序，实现数组相加功能
    arrayAdd << <1, 4 >> >(d_a, d_b, d_c);
    //利用cudaMemcpy函数将GPU端计算结果复制到CPU端
    cudaMemcpy(C, d_c, size, cudaMemcpyDeviceToHost);
    //释放所有的GPU内存（切记有借有还，不然后果自负。。。）
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    system("Pause");
    return 0;
}

CUDA编程简单来说分为3大步，首先创建在GPU中的内存空间，将有初值的变量从CPU端复制值GPU端，接着调用在GPU内运行的内核函数进行计算，最后将内核函数的运算结果复制到CPU端的变量并且释放掉所有之前创建的GPU内存（这步不能漏，不然后面会出一堆内存问题）。
以上面的程序为例，程序目的为利用并行计算将数组A与数组B相加，首先我在主机端随意定义了3个int型数组，其数组长度均为4，接着利用cudaMalloc函数对其GPU内存进行开辟，内存大小为4倍的int型，接着利用cudaMemcpy函数对数据进行传输，这里在函数的末尾“cudaMemcpyHostToDevice”为一个cudaMemcpy函数特有的标识码，即将主机端（Host）的数据传入设备端（Device）。
然后调用了内核函数arrayADD，在解释此内核函数前，首先我将引入CUDA中Grid（网格）、Block（块）、Thread（线程）的概念，每个内核函数在启动前，程序员必须对本次计算的维度进行设计，thread是最小的计算单元，多个thread组成了一个block,而多个block组成了一个grid，这里可能有同学会问，让多线程一起计算直接用相应数量的thread不就行了，为什么还要引入block和grid，增加设计的复杂度，这是因为每台电脑的显卡在硬件设计的时候都存在的一定的物理极限，简单来说一个block中的thread不是无限的，举个例子，本人目前使用的本本中，显卡型号为Nvidia Geforce 755M，它的最大线程数为1024，超过这个值程序就无法运行，故在大计算时需要合理设计block中thread的数量，这在后续的实践中我会继续提到，关于Grid（网格）、Block（块）、Thread（线程）这三者更具体的关系可参见此博客[CUDA 的Threading：Block 和Grid 的设定与Warp](http://www.360doc.com/content/13/0819/19/12278894_308350397.shtml)

由于本例数据较少，在计算任务的分配上直接利用了4个thread和一个block，如此设计在于每个thread承担一次A[i]+B[i]的任务。最后再利用cudaMemcpy函数将运算结果d_c传送给数组C，一个简单的CUDA小程序就完成了，由于本例的目的在于初步尝试CUDA编程，一些细节小问题望大家不要深究~

CUDA中线程索引threadIdx和块索引blockIdx

在上文的内核函数中，第一行中threadIdx.x为线程的索引，它的目的在于让GPU可以准确的找到目的操作数用对应的线程进行计算工作，是CUDA编程中和重要的一部分。

本人在学习CUDA的过程中，一开始经常被线程索引threadIdx和块索引blockIdx搞的崩溃，现将我在学习中总结一些的小经验进行分享：

（1）CUDA中kernel内的变量和常量

在CUDA中每个块中线程的索引变量threadIdx分为X方向上的索引threadIdx.x以及Y方向上的索引threadIdx.y，块索引变量blockIdx分为X方向上的索引blockIdx.x以及Y方向上的索引block.y，通过这两个索引我们可以准确定位到某个线程。

而kernel中的常量为blockDim以及gridDim，它们分别代表了没block的大小以及grid中block的分布。

（2）创建内核程序中block和thread

dim3 blocksPerGrid(N1,M1);//在grid中有N1*M1个block
dim3 threadsPerBlock(N2,M2);//每个block中有N2*M2个thread

（3）创建索引

int threadIndex = threadIdx.x + threadIdx.y * blockDim.x;
int blockIndex = blockIdx.x + blockIdx.y * gridDim.x;
int index = threadIndex + blockIndex * blockDim.x * blockDim.y;

threadIndex为单个block中thread的索引，blockIndex为block在整个grid中的索引，index即为某个thread在整个grid中的索引。

结束

这是我的第一篇CUDA知乎文，写的可能有些粗陋，后续会继续上传我的一些CUDA程序，希望大家多多指教~