2023 CUDA线上春训营笔记 向量相加例程归纳

GPU中向量相加的步骤，主要分为以下几个步骤：

1. 将计算机内存中的数据搬运到显存中
2. 对于N维的向量，申请N个线程，每个线程对应向量的一个维度。
3. 每个线程分别计算每隔维度的和。

#include <math.h>
#include <stdio.h>

void __global__ add(const double *x, const double *y, double *z, int count)
{
    // 获取全局索引
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    if( n < count)
    {
        //向量中的每个元素相加
        z[n] = x[n] + y[n];
    }
}

// 检查执行结果
void check(const double *z, const int N)
{
    bool error = false;
    for (int n = 0; n < N; ++n)
    {
        if (fabs(z[n] - 3) > (1.0e-10))
        {
            error = true;
        }
    }
    printf("%s\n", error ? "Errors" : "Pass");
}


int main(void)
{
    const int N = 1000;
    const int M = sizeof(double) * N;
    //申请CPU内存
    double *h_x = (double*) malloc(M);
    double *h_y = (double*) malloc(M);
    double *h_z = (double*) malloc(M);

    for (int n = 0; n < N; ++n)
    {
        h_x[n] = 1;
        h_y[n] = 2;
    }

    // 申请GPU内存
    double *d_x, *d_y, *d_z;
    cudaMalloc((void **)&d_x, M);
    cudaMalloc((void **)&d_y, M);
    cudaMalloc((void **)&d_z, M);
    cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);

    const int block_size = 128;
    const int grid_size = (N + block_size - 1) / block_size; //这句话是确保grid_size大于N
    // grid_size 和 block_size 分别代表了本次 kernel 启动对应的 block 数量和每个 block 中 thread 的数量，所以显然两者都要大于 0。
    add<<<grid_size, block_size>>>(d_x, d_y, d_z, N);

    cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
    check(h_z, N);

    //释放CPU内存
    free(h_x);
    free(h_y);
    free(h_z);
    //释放GPU内存
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    return 0;
}