CUDA编程：其三、CUDA向量加法

为什么CUDA并行计算特别适用于向量加法呢？因为计算向量加法的时候各个元素之间互相独立，互不影响。假如所我要计算两个1024长度的向量相加，那么就可以启动1024个线程，每个线程分别计算一个位置上的值。

一、重要函数介绍
（1）cudaMalloc
这个函数负责在GPU申请内存，对应CPU之下的malloc函数。
函数声明如下：

cudaError_t cudaMalloc (void **devPtr, size_t  size ); 

第一个参数表示传入的用来记录数据的指针的地址（有点拗口），第二次参数很好理解，表示申请的GPU空间的大小。
先看一个例子：

int *devPtr = nullptr;
    
// 分配GPU内存
cudaMalloc((void**)&devPtr, 100 * sizeof(int));

这表示在GPU上开辟了一个长度100的int数组。为什么第一个参数要定义成指针的指针呢？
首先回忆一下CPU下面的malloc函数：

void * malloc(size_t  size)

因为返回的是开辟的空间的首地址，调用方可以拿到这个首地址。但是cudaMalloc返回值是void，是通过出参的方式拿到这个地址。在C++里面可以使用引用变量作为入参很好地解决这个问题，但是cuda接口更偏向于c语言，只能使用指针的指针作为出参。
简单地说，第一个参数是形参，函数里面赋值的是局部变量，外界变量不会被改变。所谓指针，就是存储地址的变量，他本身也是有地址的，把devPtr指针的地址传进去，就可以改变函数外部指针的值。
很别扭，C++工程师估计得反应一会，c工程师估计很好理解。

（2）cudaMemcpy
该函数是负责在CPU和GPU之间传递数据的。
函数声明如下：

//表示从CPU到GPU拷贝数据
cudaMemcpy(d_A,h_A,nBytes,cudaMemcpyHostToDevice);

//表示从GPU到CPU拷贝数据
cudaMemcpy(h_A,d_A,nBytes,cudaMemcpyDeviceToHost);

一般的流程是，先在CPU处生成数据，之后拷贝给GPU；GPU处理好后拷贝回给CPU。

（3）cudaFree
不多说了，释放指针而已。

二、向量加法的不同实现
（1）一个block多个thread。
假如向量长度是n，一个block中线程个数是m，相对于每个线程要计算n/m个数字。看代码：

#include <iostream>
#define  N 10

__global__ void addVector(int *a, int *b, int *c, int n)
{
    int threadId = threadIdx.x;
    int step = blockDim.x;
    while(threadId < n){
        c[threadId] = a[threadId]+b[threadId];
        threadId += step;
    }
}

int main() {
    int a[N], b[N], c[N];
    int *gpu_a, *gpu_b, *gpu_c;
    for(int i=0; i<N; ++i) // 为数组a、b赋值
    {
        a[i] = i;
        b[i] = i * i;
    }
    cudaMalloc(&gpu_a, sizeof(int) * N);
    cudaMemcpy(gpu_a, a, sizeof(int) * N, cudaMemcpyHostToDevice);

    cudaMalloc(&gpu_b, sizeof(int) * N);
    cudaMemcpy(gpu_b, b, sizeof(int) * N, cudaMemcpyHostToDevice);

    cudaMalloc(&gpu_c, sizeof(int) * N);
    cudaMemcpy(gpu_c, c, sizeof(int) * N, cudaMemcpyHostToDevice);

    addVector<<<1, 3>>>(gpu_a, gpu_b, gpu_c, N);
    cudaMemcpy(c, gpu_c, sizeof(int) * N, cudaMemcpyDeviceToHost);

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

    cudaFree(gpu_a);
    cudaFree(gpu_b);
    cudaFree(gpu_c);

    return 0;
}

线程ID的作用就体现出来了，通过线程id，可以划分出来每个线程处理的区间。

（2）多个block，一个block下面多个线程：

#include <iostream>
#define  N 10

__global__ void addVector(int *a, int *b, int *c, int n)
{
    int threadId = blockIdx.x* blockDim.x + threadIdx.x;
    int step = gridDim.x * blockDim.x; //线程数量
    while(threadId < n){
        c[threadId] = a[threadId]+b[threadId];
        threadId += step;
    }
}

int main() {
    int a[N], b[N], c[N];
    int *gpu_a, *gpu_b, *gpu_c;
    for(int i=0; i<N; ++i) // 为数组a、b赋值
    {
        a[i] = i;
        b[i] = i * i;
    }
    cudaMalloc(&gpu_a, sizeof(int) * N);
    cudaMemcpy(gpu_a, a, sizeof(int) * N, cudaMemcpyHostToDevice);

    cudaMalloc(&gpu_b, sizeof(int) * N);
    cudaMemcpy(gpu_b, b, sizeof(int) * N, cudaMemcpyHostToDevice);

    cudaMalloc(&gpu_c, sizeof(int) * N);
    cudaMemcpy(gpu_c, c, sizeof(int) * N, cudaMemcpyHostToDevice);

    addVector<<<2, 2>>>(gpu_a, gpu_b, gpu_c, N);
    cudaMemcpy(c, gpu_c, sizeof(int) * N, cudaMemcpyDeviceToHost);

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

    cudaFree(gpu_a);
    cudaFree(gpu_b);
    cudaFree(gpu_c);

    return 0;
}