使用 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;
}