使用 CUDA 对两个数组求和

这个例子说明了如何创建一个简单的程序，它将两个 int 数组与 CUDA 相加。

CUDA 程序是异构的，由在 CPU 和 GPU 上运行的部分组成。

利用 CUDA 的程序的主要部分类似于 CPU 程序，包括

• 将在 GPU 上使用的数据的内存分配
• 数据从主机内存复制到 GPU 内存
• 调用内核函数来处理数据
• 将结果复制到 CPU 内存

要分配设备内存，我们使用 cudaMalloc 功能。要在设备和主机之间复制数据，可以使用 cudaMemcpy 功能。cudaMemcpy 的最后一个参数指定了复制操作的方向。有 5 种可能的类型：

• cudaMemcpyHostToHost - 主持人 - >主持人
• cudaMemcpyHostToDevice - 主机 - >设备
• cudaMemcpyDeviceToHost - 设备 - >主机
• cudaMemcpyDeviceToDevice - 设备 - >设备
• cudaMemcpyDefault - 基于默认的统一虚拟地址空间

接下来调用内核函数。三个 V 形之间的信息是执行配置，它指示有多少设备线程并行执行内核。第一个数字（示例中为 2）指定块数，第二个（示例中为 (size + 1) / 2） - 块中的线程数。请注意，在此示例中，我们将大小添加 1，以便我们请求一个额外的线程，而不是让一个线程负责两个元素。

由于内核调用是异步函数，因此调用 cudaDeviceSynchronize 以等待执行完成。将结果数组复制到主机内存，并使用 cudaFree 释放设备上分配的所有内存。

要将函数定义为内核，请使用 __global__ 声明说明符。每个线程都将调用此函数。如果我们希望每个线程处理结果数组的元素，那么我们需要一种区分和识别每个线程的方法。CUDA 定义了变量 blockDim，blockIdx 和 threadIdx。预定义变量 blockDim 包含内核启动的第二个执行配置参数中指定的每个线程块的维度。预定义变量 threadIdx 和 blockIdx 分别包含其线程块内的线程索引和网格内的线程块。请注意，由于我们可能要求比阵列中的元素多一个线程，因此我们需要传入 size 以确保我们不会访问数组的末尾。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

__global__ void addKernel(int* c, const int* a, const int* b, int size) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < size) {
        c[i] = a[i] + b[i];
    }
}

// Helper function for using CUDA to add vectors in parallel.
void addWithCuda(int* c, const int* a, const int* b, int size) {
    int* dev_a = nullptr;
    int* dev_b = nullptr;
    int* dev_c = nullptr;

    // Allocate GPU buffers for three vectors (two input, one output)
    cudaMalloc((void**)&dev_c, size * sizeof(int));
    cudaMalloc((void**)&dev_a, size * sizeof(int));
    cudaMalloc((void**)&dev_b, size * sizeof(int));

    // Copy input vectors from host memory to GPU buffers.
    cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);

    // Launch a kernel on the GPU with one thread for each element.
    // 2 is number of computational blocks and (size + 1) / 2 is a number of threads in a block
    addKernel<<<2, (size + 1) / 2>>>(dev_c, dev_a, dev_b, size);
    
    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaDeviceSynchronize();

    // Copy output vector from GPU buffer to host memory.
    cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);
}

int main(int argc, char** argv) {
    const int arraySize = 5;
    const int a[arraySize] = {  1,  2,  3,  4,  5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };

    addWithCuda(c, a, b, arraySize);

    printf("{1, 2, 3, 4, 5} + {10, 20, 30, 40, 50} = {%d, %d, %d, %d, %d}\n", c[0], c[1], c[2], c[3], c[4]);

    cudaDeviceReset();

    return 0;
}