1. Vector Add

本篇讲述如何使用CUDA编程并行计算两个向量的加法。

并行结构

这是CUDA编程的并行结构。
Grid是最高层级的单位，由多个Block(线程块)组成；
每个Block由许多Thread(线程)组成，同一个Grid内的所有Block共享相同的全局内存空间；
而Thread是CUDA编程中最小的执行单元。
由此可见，一个Grid与传统的进程的概念有所相似(进程内资源共享)，而类似的，Block与Thread都对应传统的线程，但CUDA编程中的Thread是最小的执行单元，更加轻量级。

CUDA 概念	传统概念	相似点	关键区别
Grid	Process	都表示一个独立的计算任务	Grid 是 GPU 上的任务，依赖主机进程；Process 是操作系统中的独立实体
Block	Thread	都表示并行执行单元	Block 是 GPU 上的线程组，共享快速内存；Thread 是 CPU 上的执行单元
Thread	Thread	都执行代码逻辑	CUDA 线程是轻量级的（由 GPU 硬件调度）；传统线程由操作系统调度

回过头看核函数的调用，我们提到调用核函数时需要使用func<<<X, Y>>>()的语法，实际上这里的X指定的是每个Grid的维度，即每个Grid内Block是如何分布的，Y指定的是每个Block的维度，即每个Block内Thread是如何分布的。
注意<<<X, Y>>>不能指定Grid的数量。
X和Y可以是整型或dim3。

#include <iostream>

__device__ int cnt1 = 0, cnt2 = 0; // GPU上的全局变量

__global__ void kernal(int type) {
    if (type == 0) {
        printf("int parameter <<<2, 1>>>, grid.shape=(2,), block.shape(1,)\tcnt1=%d\n", ++cnt1);
    }
    else {
        printf("dim3 parameter <<<dim3(2, 2), dim3(3, 3)>>>, grid.shape=(2, 2),"
                                "block.shape=(3, 3)\tcnt2=%d\n", ++cnt2);
    }
}

/*
__global__ void kernal(int type， int& cnt, const char* str) {
    printf("%s\tcnt%d=%d\n", str, type + 1, ++cnt);
}
*/

int main() {
    kernal<<<2, 1>>>(0);
    dim3 gridshape(2, 2);
    dim3 blockshape(3, 3);
    kernal<<<gridshape, blockshape>>>(1);
    cudaDeviceSynchronize();
    return 0;
}

我们之后会继续拿该程序来作为例子。这里不展示程序输出，建议自己运行一下。
这里先看一下这里展示的两种核函数调用方式，一种直接用整型参数设置一维的Grid和Block，也可以用dim3结构体设置一维、二维或三维的Grid和Block，代码中展示的是二维，一维和三维也是类似的，例如dim3(5)、dim3(3, 2, 4)。
代码中涉及的其他问题暂时保留。

内存分配

如何在GPU上动态分配内存？
cudaError_t cudaMalloc(void **devPtr, size_t size);
cudaMalloc()函数用于在GPU上开辟一段内存空间。简单看一下形参，devPtr是一个二级指针，这样可以修改传入的实参指针，使其指向开辟的显存空间，size自然是数组大小。
如何把CPU端的数组复制到GPU内存中去？
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);
dst和src常识，前者目的(destination)地址，后者源(source)地址，count即size。
kind为复制方向，cudaMemcpyKind是一个枚举类型，主要包含以下五种取值：

取值	含义	源地址类型	目标地址类型
cudaMemcpyHostToDevice	从主机（CPU）到设备（GPU）	主机指针	设备指针
cudaMemcpyDeviceToHost	从设备（GPU）到主机（CPU）	设备指针	主机指针
cudaMemcpyDeviceToDevice	设备到设备（GPU 内部或多 GPU 之间）	设备指针	设备指针
cudaMemcpyHostToHost	主机到主机（CPU 内部）	主机指针	主机指针
cudaMemcpyDefault	统一内存（Unified Memory）模式下的自动检测方向	统一内存指针	统一内存指针

我们想把CPU端的数组复制到GPU端，就用cudaMemcpyHostToDevice。
另外注意cudaMalloc的devPtr是二级指针，因此需要传入指针的地址，而cudaMemcpy的dst和src是一级指针，传入指针的值(指向的地址)即可。

向量加法

#include <iostream>
#include <cuda_runtime.h>
#include <numeric>

__device__ int deviceArr1[100]; // GPU上的静态全局数组

__global__ void kernal(int* deviceArr2, int size) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < size)
        deviceArr1[idx] += deviceArr2[idx];
}

__global__ void init(int size) {
    for (int i = 0; i < size; ++i) {
        deviceArr1[i] = size - i;
    }
}

__global__ void display(int* deviceArr2, int size) {
    printf("deviceArr1\n");
    for (int i = 0; i < size; ++i) {
        printf("%-3d ", deviceArr1[i]);
        if (i % 10 == 9) printf("\n");
    }
    printf("deviceArr2\n");
    for (int i = 0; i < size; ++i) {
        printf("%-3d ", deviceArr2[i]);
        if (i % 10 == 9) printf("\n");
    }
}

int main() {
    const int N = 100;
    int *deviceArr2; // GPU上的动态全局数组
    int hostArr[N];
    std::iota(hostArr, hostArr + N, 1);

    cudaError_t err = cudaMalloc(&deviceArr2, 1000);
    // 在GPU全局内存中动态分配一段内存空间
    if (err != cudaSuccess) {
        // 后面为了方便 不再做检查
        printf("Failed: %s\n", cudaGetErrorString(err));
        return 1;
    }
    init<<<1, 1>>>(N); // 不能直接在主机端操作GPU的内存(主机端需要通过CUDA API)
    cudaDeviceSynchronize();
    cudaMemcpy(deviceArr2, hostArr, sizeof(hostArr), cudaMemcpyHostToDevice);
    // 这里选择用一个线程处理一个数组元素的映射方式 因此总线程数要>=数组元素数
    // 换句话说 需要有一个映射的方式决定每个线程处理哪些元素 这里直接用一个线程对应一个元素
    int blockshape = 32;
    int gridshape = (N - 1) / blockshape + 1;
    // ceil(x / y) = floor((x - 1) / y) + 1 上取整保证总线程数>=数组元素数
    kernal<<<gridshape, blockshape>>>(deviceArr2, N);
    cudaDeviceSynchronize();
    display<<<1, 1>>>(deviceArr2, N); 
    cudaDeviceSynchronize();

    cudaFree(deviceArr2);
    return 0;
}

\(deviceArr1 \leftarrow deviceArr1 + deviceArr2\)，其中\(deviceArr1[i] = 100 - i\)，使用静态数组；\(deviceArr2[i] = i + 1\)，使用动态内存分配。