CUDA指南-CUDA编程基础

CUDA编程基础是开始利用GPU进行并行计算的起点。以下是一些入门步骤和概念：

Hello World：第一个CUDA程序

编写CUDA核函数：
创建一个简单的核函数，例如一个向量加法操作。核函数用 global 修饰，表示它将在GPU上执行。

cuda

CUDA核函数定义

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

主机代码：

编写主机代码来调用核函数，并在GPU和CPU之间传输数据。

int main() {
    int n = 256;
    int *h_a, *h_b, *h_c;  // 主机上的向量
    int *d_a, *d_b, *d_c;  // 设备上的向量

    // 分配主机内存
    h_a = (int *)malloc(n * sizeof(int));
    h_b = (int *)malloc(n * sizeof(int));
    h_c = (int *)malloc(n * sizeof(int));

    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_a[i] = i;
        h_b[i] = 2 * 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);

    // 计算线程块大小和网格大小
    int blockSize = 256;
    int numBlocks = (n + blockSize - 1) / blockSize;

    // 调用核函数
    add<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);

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

    // 验证结果
    for (int i = 0; i < n; i++) {
        if (h_c[i] != h_a[i] + h_b[i]) {
            printf("Error: %d + %d != %d\n", h_a[i], h_b[i], h_c[i]);
            break;
        }
    }

    // 清理
    free(h_a);
    free(h_b);
    free(h_c);
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    return 0;
}

在CUDA编程中，确定线程块大小（blockSize）和网格大小（numBlocks）是并行计算中的关键步骤。这两者决定了你的核函数如何在GPU上执行。

线程块大小（blockSize）

线程块大小是指每个线程块中包含的线程数量。这个值通常基于以下考虑确定：

GPU架构：不同的GPU架构支持的每个线程块的最大线程数不同。例如，早期的GPU可能支持每个线程块最多1024个线程，而现代GPU可能支持更多。
内存需求：每个线程块内的线程共享共享内存，如果共享内存需求过高，可能需要减小线程块大小。
性能优化：不同的线程块大小可能会对性能产生影响，通常需要通过实验来确定最优大小。

网格大小（numBlocks）

网格大小是指执行核函数时使用的线程块的数量。它是根据总的计算任务量和每个线程块的线程数来确定的：

总任务量：n 是需要处理的总数据项数，例如向量的长度。
每个线程的任务量：通常，每个线程处理一个数据项。因此，如果blockSize是256，那么每个线程块可以处理256个数据项。

计算方法

计算网格大小的公式是：

这里向上取整，确保所有数据项都能被处理。即使n不能被blockSize整除，这个计算也能确保覆盖所有的n项数据。

在代码中，这个计算通常使用以下形式：

int numBlocks = (n + blockSize - 1) / blockSize;

这个表达式首先将n与blockSize - 1相加，确保结果至少与n一样大（因为如果n正好是blockSize的倍数，加1后再做除法会导致numBlocks比实际需要的少一个），然后除以blockSize来得到需要的线程块数量。

示例
假设有一个向量n = 1000，你选择的线程块大小是blockSize = 256：

如果直接计算1000 / 256，结果是3.92，这意味着有3.92个线程块可以完全填满，但仍有部分数据未处理。
使用上述公式，(1000 + 256 - 1) / 256 = 1004 / 256 = 3.94，向上取整得到4，这意味着你需要4个线程块来处理所有的1000个数据项。

内存管理

CUDA提供了几种不同的内存类型，每种都有其用途和性能特点：

全局内存：GPU上的主存储区，所有线程都可以访问。访问速度相对较慢。
共享内存：高速片上内存，只能被同一个线程块内的线程访问。适合线程间的数据共享。
常量内存：对于所有线程只读，可以被所有线程访问。适用于不变数据，如循环不变系数。

线程和块

在CUDA中，线程和块是实现并行计算的基本单元：

线程：执行计算的最小单元。在核函数中，每个线程执行相同的代码，但可以处理不同的数据。

线程块：一组协同工作的线程集合。线程块内的线程可以通过共享内存和同步操作来协作。

线程块的维度可以用 dim3 类型定义，例如 dim3 block(4, 4, 4);。
线程块的大小通常选择为2的幂，如256、512等，以适应GPU的执行模型。
网格：由多个线程块组成，用于在更大的范围内实现并行计算。

要创建和管理线程和块，你需要指定核函数的执行配置，即定义每个线程块的线程数和网格的线程块数。例如：

// 执行配置：定义网格和块的维度
int blockSize = 256; // 每个块256个线程
int numBlocks = (n + blockSize - 1) / blockSize; // 计算需要的线程块数量
add<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n); // 调用核函数

在编写CUDA程序时，你需要考虑如何分配数据到线程和线程块，以及如何同步线程执行，以实现高效的并行计算。