CUDA C编程：第一个程序向量相加

我的电脑没有装CUDA，所以使用租了带GPU的云服务器，然后使用vscode SSH远程连接云服务器。云GPU使用的是智星云，0.8元/h。

智星云

可以使用nvcc --version查看系统中安装的CUDA版本。

然后写第一个CUDA程序，两个向量相加结果给到第三个向量

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

#define CHECK(call) \
{ \
    const cudaError_t error = call; \
    if (error != cudaSuccess) { \
        std::cerr << "Error: " << __FILE__ << ", line " << __LINE__ << ": " \
                  << cudaGetErrorString(error) << std::endl; \
        exit(1); \
    } \
}

__global__ void addArrays(const int *A, const int *B, int *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N)
        C[idx] = A[idx] + B[idx];
}

int main() {
    const int N = 100; // 数组大小
    int A[N], B[N], C[N];

    // 初始化数组A和B
    for(int i = 0; i < N; ++i) {
        A[i] = i;
        B[i] = i * 2;
    }
    int *d_A, *d_B, *d_C;
    // 分配GPU内存
    CHECK(cudaMalloc((void**)&d_A, N * sizeof(int)));
    CHECK(cudaMalloc((void**)&d_B, N * sizeof(int)));
    CHECK(cudaMalloc((void**)&d_C, N * sizeof(int)));
    // 将数据从主机复制到设备
    CHECK(cudaMemcpy(d_A, A, N * sizeof(int), cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_B, B, N * sizeof(int), cudaMemcpyHostToDevice));
    // 调用核函数
    addArrays<<<10, 10>>>(d_A, d_B, d_C, N);
    // 同步以确保核函数执行完成
    cudaDeviceSynchronize();
    // 将结果从设备复制回主机
    CHECK(cudaMemcpy(C, d_C, N * sizeof(int), cudaMemcpyDeviceToHost));
    // 释放GPU内存
    CHECK(cudaFree(d_A));
    CHECK(cudaFree(d_B));
    CHECK(cudaFree(d_C));
    // 输出结果
    for(int i = 0; i < N; ++i)
        std::cout << C[i] << " "; // 应该输出 i + i*2
    return 0;
}

nvcc -o add add.cu编译程序

./add运行程序

程序说明

#include <cuda_runtime.h>

引入cuda运行时环境

#define CHECK(call) \
{ \
    const cudaError_t error = call; \
    if (error != cudaSuccess) { \
        std::cerr << "Error: " << __FILE__ << ", line " << __LINE__ << ": " \
                  << cudaGetErrorString(error) << std::endl; \
        exit(1); \
    } \
}

用来提供CUDA报错信息的宏，用CHECK宏嵌套每一个将要调用的函数，便于调试。

#define CHECK(call) 定义了一个名为CHECK的宏，它接受一个参数call，这个参数是想检查的CUDA API调用。接下来的花括号 { ... } 包围了宏展开后将要执行的代码块。

const cudaError_t error=call；执行传入的CUDA API调用（即call），并将其返回的错误状态保存在变量error中。

if(error!=cudaSuccess){...}：检查error是否等于cudaSuccess，这是CUDA中表示操作成功的常量。如果不等于(即操作失败)，则执行大括号内的错误处理代码。

std::cerr<< "Error: " <<__FILE__<<", line "<<__LINE__<<": "<<cudaGetErrorString(error)<< std::endl; 这行代码打印错误信息到标准错误输出。包括了出错的文件名（由__FILE__宏提供）、行号（由__LINE__宏提供），以及通过cudaGetErrorString(error)获取的错误描述字符串。exit(1); 如果确实发生了错误，程序会调用exit(1)立即终止，返回码1通常表示异常终止。

__global__ void addArrays(const int *A, const int *B, int *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N)
        C[idx] = A[idx] + B[idx];
}

__global__ 是一个关键字，用于声明一个在GPU上执行的函数，也称为全局函数或内核函数。这些函数由主机（CPU）调用，但在设备（GPU）上的多个线程并行执行。

void addArrays(const int *A,const int *B,int *C, int N）定义了内核函数addArrays。

const int *A 和 const int *B指向输入数组A和B的指针，在内核中只读。

int *C输出数组C的指针，存放A和B对应元素的和。

N:需要相加的元素个数。

int idx = blockIdx.x * blockDim.x + threadIdx.x; 计算当前线程的全局索引 idx。

这里是CUDA线程组织方式的一个体现：

blockIdx.x 是当前线程所在的块(block)在网格(grid)中的x轴索引。

blockDim.x 是每个块中线程的数量(块的尺寸)在x轴方向。

threadIdx.x 是当前线程在块内的x轴索引。通过这样的计算，每个线程都能知道自己在整个计算任务中的唯一位置，从而决定应该处理哪个数组元素。

if (idx < N) 是一个边界检查，确保线程不会访问超过数组界限。因为CUDA会为整个网格启动比实际需要更多的线程以充分利用硬件资源，所以这种检查是必要的。

C[idx] = A[idx] + B[idx]; 如果索引idx在有效范围内，这个语句就执行数组A和B中相应位置的元素相加，并将结果存储到数组C的相同位置。

（这个地方还是没怎么看懂）。

cudaMalloc((void**)&d_A, N * sizeof(int))

给设备分配N个int类型的内存，使用指针变量d_A指示。

cudaMemcpy(d_A, A, N * sizeof(int), cudaMemcpyHostToDevice)

内存拷贝，从Host拷贝到Device。A数组赋值给d_A数组。