CUDA基础的基础教程一、CUDA入门

辣鸡博主又开新坑……
CUDA是NVIDA家的GPU编程套件，他通过提供基本的底层执行环境和并行编程库，为GPU编程提供了便利。
介绍到此为止，那么CUDA的程序是怎样工作的呢？
首先，我们用GPU来做什么？当然不是用来输出一个“hello world”，我们用CUDA进行图形渲染，大数据计算，科学计算，深度学习等数据量巨大的运算。虽然我们的CPU是足够强大的，然而一个CPU的能力有限，而GPU用他的高并发来解决这些问题，相当于是将一个问题不断地分解成子问题，完成巨大计算任务。

当然博主本身是CV出身，有关图形学的内容基本不会涉及……
除此之外，我们会讲讲高级的API，基础部分算是一个简单的入门

1.我的第一个CUDA程序

我们说了这么多，就从一个简单的问题入手吧。我们现在有两个矢量（或者说是数组$A_i,B_i$ 我们现在想求$C_i = A_i + B_i$ :

void add(int *a,int *b ,int *c)
{
    for(int i = 0;i<N;i++) c[i] = a[i] + b[i] //N - > size of arrays
}

这是简单的加法操作，相信读者都可以接受，那么我们来分析一下这个问题，最占运算的是什么。很明显是中间的循环。举个例子，工人需要从一个地方搬东西到另一个地方，一个人的话就是N趟，时间就是N×一个人的时间。那么假如我们有N个人呢？我们就可以在一个人搬运所需的时间内完成十个人搬运的时间内完成这项工作。
这项任务我们来用CUDA实现，先来CUDA中交给每个创建的进程处理的函数叫做kernel function核函数，我们先来看核函数的定义：

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

首先在声明最前面是一个__global__ 声明，代表这个函数可以在GPU和CPU上调用，如果这个声明是__device__ 就代表这个函数只能在设备（GPU）上运行。
我们把完整的代码放上来再说

#include <iostream>
#include "book.h"
#define HD      HANDLE_ERROR
#define mloc cudaMalloc
#define mcpy cudaMemcpy
#define mfree cudaFree
#define d2h  cudaMemcpyDeviceToHost
#define h2d cudaMemcpyDeviceToHost    
using namespace std;
#define N 100000
__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;


    HD(mloc((void**)&dev_a,N*sizeof(int)));
    HD(mloc((void**)&dev_b,N*sizeof(int)));
    HD(mloc((void**)&dev_c,N*sizeof(int)));


    for(int i = 0;i<N;i++)
    {
        a[i] = -i;
        b[i] = i *i;
    }
    HD(mcpy(dev_a,a,N*sizeof(int),h2d));
    HD(mcpy(dev_b,b,N*sizeof(int),h2d));

    add<<<N,1>>>(dev_a,dev_b,dev_c);
    HD(mcpy(c,dev_c,N*sizeof(int),d2h));
    for(int i = 0;i<N;i++)
    {
        printf("%d + %d = %d\n",a[i],b[i],c[i]);
    }
    mfree(dev_a);
    mfree(dev_b);
    mfree(dev_c);
    return 0;

}

我们注意调用这一行的代码：

add<<<N,1>>>(dev_a,dev_b,dev_c);

三个尖括号是调用GPU操作的标志，第一个参数代表进程块的数量，第二个参数代表每个进程块内有多少个进程，上面那个函数直接调用了10000个进程块来执行这个操作，一个线程块只有一个线程，我们先不管进程块的概念，直接把这个块当成一个进程。那么就是就是10000个进程，这样的执行效率是很快的，但是随之而来的是，一个进程怎么才能知道自己要执行那个操作？
我们观察add函数的第一行，会发现blockIdx.x这个变量，这是CUDA内置的标示变量，这个就是代表是第几个进程块，相当于我们给10000个人分配任务，每个人拿到的任务就是他对应id的数组成员的和。
我们还要注意：

    HD(mloc((void**)&dev_a,N*sizeof(int)));
    HD(mloc((void**)&dev_b,N*sizeof(int)));
    HD(mloc((void**)&dev_c,N*sizeof(int)));

要知道我们在GPU上无法读取内存中的数据，只能在GPU上声明内存，用法和CPU是一样的。在最后我们同样要释放自己的申请资源：

    mfree(dev_a);
    mfree(dev_b);
    mfree(dev_c);

这就出现了另一个问题，既然内存个GPU内存不通用，就需要进行交换：

    HD(mcpy(dev_a,a,N*sizeof(int),h2d));
    HD(mcpy(dev_b,b,N*sizeof(int),h2d));

这个相对比C中的memcpy，实际上只有最后一个参数，也就是方向的区别，这个函数支持双向操作，很是方便。
于是一个CUDA版的的“hello world”就完成了，当然我们还注意到了一个地方：

    if(tid < N) c[tid] = a[tid] * b[tid];

这里为什么要检查一下呢？这就相当于一种非法检查机制，要是在GPU上发生了非法访问，可是很麻烦的，最好注意一下。
辣鸡博主喜欢写宏，请大家仔细看宏定义……
这个book.h是在CUDA_BY_Example 的官网上下载的。
大家试着运行一下吧，贼快～

2.总结：

我们可以想象CUDA是一个任务分配系统，每个线程块（相当于是工人小队）中有若干的线程（工人），通过调用核函数来让CUDA分配的工作，每个线程通过自己所处的位置来判断自己的任务是什么。

3. 相关资料：

CUDA_BY_Example :https://developer.nvidia.com/cuda-example