First class of cuda: add kernel

• thread: 最小执行单位
• block: 线程块，一组thread 的集合
• grid: 多个block 组成网格，即一次kernel 启动的全部线程
• threadIdx: 当前线程在 block 中的索引
• blockIdx: 当前线程所在 block 在 grid 中的索引

内存分成两边：

• CPU / Host / 主机内存
• GPU / Device / 设备显存

malloc 分配的是 CPU 内存：

float *h_a = static_cast<float *>(std::malloc(bytes));

这里的 h_ 通常表示 host，意思是这块内存在 CPU 这边。

cudaMalloc 分配的是 GPU 显存：

CUDA_CHECK(cudaMalloc(&d_a, bytes));
CUDA_CHECK(cudaMalloc(&d_b, bytes));
CUDA_CHECK(cudaMalloc(&d_c, bytes));

这里的 d_ 通常表示 device，意思是这块内存在 GPU 这边。

cpu 不能直接访问 d_c，需要通过 cudaMemcpy 从 GPU 复制到 CPU：

CUDA_CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));

cudaDeviceSynchronize() 的作用是让 CPU 等 GPU 当前任务执行完成。它也会捕获 kernel 运行过程中发生的错误，比如非法访问显存。

所以这两句常常一起出现：

CUDA_CHECK(cudaGetLastError());       // 检查 kernel 是否成功启动
CUDA_CHECK(cudaDeviceSynchronize());  // 等 kernel 执行完，并检查执行时错误

Full example

#include <cmath>
#include <cstdio>
#include <cstdlib>

#define CUDA_CHECK(call)                                                       \
    do {                                                                       \
        cudaError_t err = (call);                                               \
        if (err != cudaSuccess) {                                               \
            std::fprintf(stderr, "CUDA error %s:%d: %s\n", __FILE__, __LINE__, \
                         cudaGetErrorString(err));                             \
            std::exit(EXIT_FAILURE);                                            \
        }                                                                      \
    } while (0)

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

int main() {
    const int n = 1 << 20;
    const size_t bytes = n * sizeof(float);

    float *h_a = static_cast<float *>(std::malloc(bytes));
    float *h_b = static_cast<float *>(std::malloc(bytes));
    float *h_c = static_cast<float *>(std::malloc(bytes));
    if (!h_a || !h_b || !h_c) {
        std::fprintf(stderr, "host malloc failed\n");
        return EXIT_FAILURE;
    }

    for (int i = 0; i < n; ++i) {
        h_a[i] = static_cast<float>(i) * 0.5f;
        h_b[i] = static_cast<float>(i) * 2.0f;
    }

    float *d_a = nullptr;
    float *d_b = nullptr;
    float *d_c = nullptr;
    CUDA_CHECK(cudaMalloc(&d_a, bytes));
    CUDA_CHECK(cudaMalloc(&d_b, bytes));
    CUDA_CHECK(cudaMalloc(&d_c, bytes));

    CUDA_CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));

    const int threadsPerBlock = 256;
    const int blocks = (n + threadsPerBlock - 1) / threadsPerBlock;
    vectorAddKernel<<<blocks, threadsPerBlock>>>(d_a, d_b, d_c, n);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());

    CUDA_CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));

    int errors = 0;
    for (int i = 0; i < n; ++i) {
        float expected = h_a[i] + h_b[i];
        if (std::fabs(h_c[i] - expected) > 1e-5f) {
            if (errors < 10) {
                std::fprintf(stderr, "mismatch at %d: got %f, expected %f\n", i,
                             h_c[i], expected);
            }
            ++errors;
        }
    }

    if (errors == 0) {
        std::printf("PASS: checked %d elements\n", n);
    } else {
        std::printf("FAIL: %d mismatches\n", errors);
    }

    CUDA_CHECK(cudaFree(d_a));
    CUDA_CHECK(cudaFree(d_b));
    CUDA_CHECK(cudaFree(d_c));
    std::free(h_a);
    std::free(h_b);
    std::free(h_c);

    return errors == 0 ? EXIT_SUCCESS : EXIT_FAILURE;
}