CUDA C/C++编程

在C++中直接使用CUDA进行GPU编程是CUDA最原生的支持方式。
CUDA C/C++是一种扩展了C/C++语言的编程模型，它允许开发者编写专门在NVIDIA GPU上执行的代码。
以下是一些基本的CUDA C/C++编程概念和常用指令：

CUDA C/C++ 基本概念

1、Kernel 函数

定义：Kernel 函数是在 GPU 上执行的函数，由主机代码调用。
声明：Kernel 函数必须使用 global 关键字进行声明。
调用：Kernel 函数由主机代码调用，通常使用 <<< >>> 语法指定线程块和网格的维度。
示例：

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

2、线程、线程块和网格

线程：CUDA 中的线程是最小的并行执行单位。
线程块：线程被组织成线程块，每个线程块内的线程可以协同执行。
网格：线程块进一步组织成网格，网格中的所有线程块一起构成一个执行单位。
示例：

// 设置线程块和网格的尺寸
dim3 blockSize(256); // 每个线程块中有256个线程
dim3 gridSize((n + blockSize.x - 1) / blockSize.x); // 网格中线程块的数量
// 调用 Kernel
add<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

3、设备和主机内存

设备内存：位于 GPU 上的内存，用于存储数据和中间结果。
主机内存：位于 CPU 上的内存，用于存储原始数据和最终结果。
数据传输：数据必须从主机复制到设备才能在 GPU 上处理，处理完成后数据再复制回主机。
示例：

// 分配设备内存
cudaMalloc((void**)&d_a, n * sizeof(int));
cudaMalloc((void**)&d_b, n * sizeof(int));
cudaMalloc((void**)&d_c, n * sizeof(int));

// 将数据从主机复制到设备
cudaMemcpy(d_a, h_a, n * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, n * sizeof(int), cudaMemcpyHostToDevice);

// 将结果从设备复制回主机
cudaMemcpy(h_c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost);

// 释放设备内存
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);

4. 同步

线程同步：在同一个线程块内部，可以使用 __syncthreads() 函数来同步线程。
线程块同步：不同线程块之间的同步通常通过正确的网格和线程块配置来实现，或者在 Kernel 函数外通过 CUDA 事件或流来实现。

5. 内存层次结构

全局内存：位于设备上，所有线程都可以访问，但访问速度较慢。
共享内存：位于设备上，同一线程块内的线程可以快速访问。
寄存器：位于每个线程中，访问速度最快，但数量有限。
常量内存：只读，所有线程都可以访问，用于存储不变的数据。
纹理内存：优化了2D访问模式的内存。

6. 错误处理

检查错误：每次调用 CUDA API 之后，检查是否有错误发生。
示例：

cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("CUDA error: %s\n", cudaGetErrorString(err));
}

常用指令示例

假设你有一个简单的向量加法操作，下面是如何使用CUDA C/C++来实现它的代码：
示例：

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

// 定义一个将在GPU上执行的Kernel函数
__global__ void add(int *a, int *b, int *c, int n) {
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < n) {
        c[index] = a[index] + b[index];
    }
}

int main() {
    const int N = 1024; // 向量大小
    int h_a[N], h_b[N], h_c[N]; // 主机上的向量
    int *d_a, *d_b, *d_c; // 设备上的向量

    // 初始化主机向量
    for (int i = 0; i < N; i++) {
        h_a[i] = i;
        h_b[i] = i;
    }

    // 分配设备内存
    cudaMalloc((void**)&d_a, N * sizeof(int));
    cudaMalloc((void**)&d_b, N * sizeof(int));
    cudaMalloc((void**)&d_c, N * sizeof(int));

    // 将数据从主机复制到设备
    cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);

    // 设置Kernel参数和调用Kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    add<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, N);

    // 将结果从设备复制回主机
    cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);

    // 清理设备内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    // 打印前几个元素验证结果
    for (int i = 0; i < 5; i++) {
        std::cout << "h_c[" << i << "] = " << h_c[i] << std::endl;
    }

    return 0;
}

这个例子展示了如何在CUDA C/C++中：

1、定义和调用Kernel函数。
2、管理设备和主机内存。
3、从主机复制数据到设备，以及反方向复制。
4、设置Kernel的网格和线程块参数。