第三章 CUDA编程入门：CUDA C/C++基础、核函数编写、线程索引计算、内存分配与数据拷贝

CUDA（Compute Unified Device Architecture）是NVIDIA推出的通用并行计算平台，其核心思想是将GPU视为一个高度并行的协处理器，通过CUDA C/C++语言扩展，让开发者能够编写在GPU上运行的并行程序。本章将围绕CUDA编程的四个核心基础模块展开：CUDA C/C++语言扩展与程序结构、核函数（Kernel）的编写规范、线程索引（Thread Index）的计算逻辑，以及GPU内存的分配与主机-设备数据拷贝。理解这些基础是后续使用cuBLAS等库和进行性能优化的前提。下图展示了CUDA编程模型的核心概念关系，包括主机与设备的分工、线程层次结构以及内存拷贝路径。

3.1 CUDA C/C++基础与程序结构

核心原理：异构计算与语言扩展

CUDA C/C++是标准C/C++语言的超集，通过添加少量关键字和运行时库函数，使开发者能够编写在GPU上执行的代码。其核心是异构计算模型：CPU作为主机（Host）负责串行逻辑、任务调度和内存管理，GPU作为设备（Device）负责大规模并行计算。一个完整的CUDA程序通常包含主机代码和设备代码两部分，设备代码通过核函数（Kernel）的形式定义，并在主机代码中通过特殊的调用语法启动。

CUDA程序的基本编译和执行流程如下：

1. 编写源代码：使用.cu扩展名保存CUDA源文件，其中可以混合主机代码（C/C++）和设备代码（CUDA扩展）。
2. 编译：使用nvcc编译器（NVIDIA CUDA Compiler）将.cu文件编译为可执行文件。nvcc会分离主机代码和设备代码，设备代码被编译为PTX（Parallel Thread Execution）或SASS（Shader Assembly）代码，主机代码则通过标准C++编译器编译。
3. 运行：程序启动后，主机代码在CPU上执行，当遇到核函数调用时，将任务加载到GPU上执行。

下表总结了CUDA C/C++中常用的关键字和内置变量：

操作流程：编写第一个CUDA程序

编写CUDA程序需要遵循以下步骤：

1. 包含头文件：#include <cuda_runtime.h> 或 #include <cuda.h>。
2. 定义核函数：使用__global__关键字声明，函数返回类型必须为void。
3. 分配设备内存：使用cudaMalloc在GPU上分配显存。
4. 拷贝数据：使用cudaMemcpy将数据从主机拷贝到设备。
5. 配置执行配置：通过<<<gridDim, blockDim>>>指定网格和块的大小。
6. 启动核函数：在主机代码中调用核函数。
7. 拷贝结果：将计算结果从设备拷贝回主机。
8. 释放内存：使用cudaFree释放设备内存。

// 核函数定义
__global__ void vectorAdd(float *A, 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() {
    int n = 1024;
    size_t bytes = n * sizeof(float);
    float *h_A, *h_B, *h_C;
    float *d_A, *d_B, *d_C;
    // 分配主机内存并初始化...
    // 分配设备内存
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);
    // 拷贝数据到设备
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
    // 配置执行配置并启动核函数
    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    vectorAdd<<<blocks, threads>>>(d_A, d_B, d_C, n);
    // 拷贝结果回主机
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
    // 释放内存
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    return 0;
}

常见误区与总结

• 误区1：忘记检查CUDA API返回值。CUDA API调用失败时返回错误码，应使用cudaGetLastError()或cudaError_t检查错误。
• 误区2：核函数中访问主机内存。核函数只能访问设备内存（显存），不能直接访问主机内存。
• 误区3：忽略执行配置的合法性。块大小（blockDim）不能超过设备限制（通常为1024），网格大小（gridDim）也有上限。

3.2 核函数编写规范

核心原理：核函数的约束与执行模型

核函数是CUDA并行计算的核心，它定义了每个线程需要执行的操作。核函数具有以下关键特性：

• 并行执行：核函数被多个线程同时执行，每个线程处理不同的数据。
• 无返回值：核函数必须返回void，结果通过参数传递（通常是指向设备内存的指针）。
• 线程同步：核函数内部可以使用__syncthreads()进行块内线程同步，但无法在全局范围内同步。
• 递归限制：核函数不支持递归调用（除非使用动态并行，但复杂度较高）。

核函数的执行模型基于SIMT（Single Instruction, Multiple Thread）架构，即所有线程执行相同的指令，但处理不同的数据。这种模型要求开发者避免线程分支发散（Warp Divergence），即同一Warp（32个线程）内的线程执行不同路径时，会导致性能下降。

操作流程：核函数编写要点

1. 确定线程索引：根据数据规模设计线程索引映射，确保每个数据元素被一个线程处理。
2. 边界检查：当数据规模不是块大小的整数倍时，需要添加边界检查（如if (i < n)）。
3. 优化内存访问：尽量使用合并访问（Coalesced Access），即相邻线程访问相邻内存地址。
4. 减少分支发散：避免在核函数中使用依赖于线程索引的条件分支。

案例拆解：矩阵乘法核函数

矩阵乘法是GPU计算的经典案例。假设我们要计算C = A * B，其中A为M×K矩阵，B为K×N矩阵，C为M×N矩阵。最简单的核函数实现如下：

__global__ void matMul(float *A, float *B, float *C, int M, int K, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < M && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < K; k++) {
            sum += A[row * K + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

该实现虽然正确，但性能较差，因为每个线程需要多次访问全局内存。后续章节将介绍使用共享内存和分块（Tiling）技术优化矩阵乘法。

3.3 线程索引计算

核心原理：线程层次与索引映射

CUDA的线程层次结构分为三层：网格（Grid）、块（Block）和线程（Thread）。一个网格包含多个块，一个块包含多个线程。每个线程可以通过内置变量threadIdx、blockIdx、blockDim和gridDim计算自己的全局索引。

对于一维数据，全局索引计算公式为：

int tid = blockIdx.x * blockDim.x + threadIdx.x;

对于二维数据（如矩阵），通常使用二维块和网格，索引计算如下：

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int idx = row * width + col;  // 行主序存储

下表对比了一维、二维和三维索引的计算方式：

操作流程：索引计算步骤

1. 确定数据维度：根据数据形状选择线程网格的维度（1D、2D或3D）。
2. 选择块大小：块大小通常为32的倍数（如128、256），以匹配Warp大小。
3. 计算网格大小：网格大小 = (数据大小 + 块大小 - 1) / 块大小。
4. 编写索引公式：在核函数中根据维度计算全局索引。

常见误区与总结

• 误区1：块大小设置过大。块大小超过设备限制（如1024）会导致启动失败。
• 误区2：网格大小计算错误。当数据大小不能被块大小整除时，需要向上取整，并在核函数中添加边界检查。
• 误区3：忽略线程索引的起始值。CUDA的索引从0开始，与C语言数组索引一致。

3.4 内存分配与数据拷贝

核心原理：主机内存与设备内存的分离

GPU拥有独立的显存（设备内存），主机无法直接访问。因此，数据在主机和设备之间传递必须通过PCIe总线，使用cudaMemcpy函数。CUDA提供了多种内存管理函数，包括：

• cudaMalloc：在设备上分配线性内存。
• cudaFree：释放设备内存。
• cudaMemcpy：在主机和设备之间拷贝数据。
• cudaMemset：将设备内存设置为指定值。

内存拷贝的方向由cudaMemcpyKind参数指定：

操作流程：内存管理最佳实践

1. 分配内存：使用cudaMalloc分配设备内存，注意检查返回值。
2. 数据初始化：在主机上初始化数据，然后拷贝到设备。
3. 执行计算：启动核函数进行计算。
4. 结果回传：将计算结果从设备拷贝回主机。
5. 释放资源：使用cudaFree释放设备内存，避免内存泄漏。

案例拆解：内存分配与拷贝的完整流程

以下代码展示了内存分配与数据拷贝的完整流程，并加入了错误检查：

#include <cstdio>
#include <cuda_runtime.h>

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

int main() {
    const int N = 1024;
    const size_t bytes = N * sizeof(float);
    float *h_data = nullptr, *d_data = nullptr;
    // 分配主机内存
    h_data = (float*)malloc(bytes);
    for (int i = 0; i < N; i++) h_data[i] = i * 1.0f;
    // 分配设备内存
    CHECK_CUDA(cudaMalloc(&d_data, bytes));
    // 拷贝数据到设备
    CHECK_CUDA(cudaMemcpy(d_data, h_data, bytes, cudaMemcpyHostToDevice));
    // 在设备上进行计算（此处省略核函数）
    // 拷贝结果回主机
    CHECK_CUDA(cudaMemcpy(h_data, d_data, bytes, cudaMemcpyDeviceToHost));
    // 验证结果...
    // 释放内存
    CHECK_CUDA(cudaFree(d_data));
    free(h_data);
    return 0;
}

该案例中，CHECK_CUDA宏用于检查每个CUDA调用的返回值，确保错误能够被及时发现。在实际项目中，这种错误检查机制是必不可少的。

常见误区与总结

• 误区1：忘记释放设备内存。设备内存泄漏会导致显存耗尽，程序崩溃。
• 误区2：使用malloc分配设备内存。必须使用cudaMalloc，否则核函数无法访问。
• 误区3：忽略内存拷贝的同步性。cudaMemcpy是同步操作，会阻塞主机直到拷贝完成。对于异步操作，可以使用cudaMemcpyAsync和流（Stream）。

本章内容为CUDA编程奠定了坚实基础。掌握核函数编写、线程索引计算和内存管理是后续学习cuBLAS库和性能优化的关键。建议读者动手编写并运行上述示例代码，通过实践加深理解。在实际项目中，务必养成检查CUDA错误、合理设计线程层次和优化内存访问的习惯，这将直接影响程序的正确性和性能。