📚 课程目录第4章 CUDA编程入门

4. CUDA编程入门：环境搭建、核函数编写、内存管理

4.1 核心原理：从“加速卡”到“并行计算平台”

CUDA（Compute Unified Device Architecture）是英伟达提供的并行计算平台与编程模型。它的核心思想并非让CPU“闲置”，而是构建CPU（Host，主机端）与GPU（Device，设备端）协同工作的异构计算体系。CPU擅长复杂逻辑与串行控制，GPU则拥有数千个轻量级核心，适合大规模数据并行任务。

理解CUDA程序执行的关键在于网格-线程块-线程的三层抽象模型。当你调用一个核函数（Kernel）时，实际上是在GPU上启动一个网格（Grid），网格由多个线程块（Block）组成，每个块内又包含若干线程（Thread）。硬件调度器会将线程块映射到GPU的流多处理器（SM）上执行。这种分层设计使得程序能自动适应不同规格的GPU——小GPU上块串行执行，大GPU上块并行执行，无需修改代码。

核函数通过 __global__ 修饰符声明，调用时使用 <<<gridDim, blockDim>>> 语法指定网格与线程块维度。GPU 硬件自动完成线程调度，开发者只需关注数据如何划分到线程上。

4.2 环境搭建：三步准备开发环境

在数据中心服务器上搭建 CUDA 开发环境，通常需要以下三个步骤：

安装 NVIDIA 驱动：确保 GPU 能被系统识别。使用 nvidia-smi 验证驱动版本与 GPU 状态。
安装 CUDA Toolkit：包含编译器 nvcc、运行时库和开发工具。注意 Toolkit 版本需与驱动版本兼容（可通过 nvidia-smi 右上角查看支持的 CUDA 版本）。
配置环境变量：将 nvcc 和库路径加入 PATH 与 LD_LIBRARY_PATH。

# 验证驱动
nvidia-smi
# 输出示例：Driver Version: 535.104.05  CUDA Version: 12.2

# 安装 CUDA Toolkit（以 Ubuntu 为例）
wget https://developer.download.nvidia.com/compute/cuda/12.2.0/local_installers/cuda_12.2.0_535.54.03_linux.run
sudo sh cuda_12.2.0_535.54.03_linux.run

# 配置环境变量（写入 ~/.bashrc）
export PATH=/usr/local/cuda/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH

# 验证编译器
nvcc --version

版本兼容性提示：数据中心服务器常使用容器化部署。推荐直接拉取官方镜像 nvcr.io/nvidia/cuda:12.2.0-devel-ubuntu22.04，其中已包含完整开发工具链，可避免宿主机环境冲突。

4.3 核函数编写：向量加法实战

我们以经典的向量加法为例，展示核函数编写与调用的完整流程。假设需要计算 C[i] = A[i] + B[i]，其中向量长度为 N。

// vecAdd.cu
#include <stdio.h>
#include <cuda_runtime.h>

// 核函数定义：每个线程处理一个元素
__global__ void vecAddKernel(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() {
    int n = 1 << 20;  // 1,048,576 个元素
    size_t bytes = n * sizeof(float);

    // 1. 分配主机端内存并初始化
    float *h_A, *h_B, *h_C;
    h_A = (float*)malloc(bytes);
    h_B = (float*)malloc(bytes);
    h_C = (float*)malloc(bytes);
    for (int i = 0; i < n; i++) { h_A[i] = 1.0f; h_B[i] = 2.0f; }

    // 2. 分配设备端内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);

    // 3. 将输入数据从主机拷贝到设备
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

    // 4. 配置核函数执行参数并启动
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vecAddKernel<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);

    // 5. 将结果拷贝回主机
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

    // 6. 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %.1f\n", i, h_C[i]);
    }

    // 7. 释放内存
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);
    return 0;
}

编译与运行命令：

nvcc -o vecAdd vecAdd.cu
./vecAdd

代码中 blockIdx.x * blockDim.x + threadIdx.x 是 CUDA 编程中最常用的全局线程索引计算公式。边界检查 if (i < n) 确保当向量长度不是线程块大小的整数倍时，不会发生越界访问。

4.4 内存管理：显存分配与数据搬运

CUDA 内存管理是性能优化的核心环节。下表总结了常用内存操作 API 及其用途：

API 函数	功能说明	注意事项
`cudaMalloc()`	在设备端分配全局内存	分配粒度建议对齐到 256 字节
`cudaMemcpy()`	主机与设备间数据拷贝	同步操作，大数据量时耗时显著
`cudaMemset()`	设备内存初始化	类似 `memset()`
`cudaFree()`	释放设备内存	必须与 `cudaMalloc` 配对
`cudaHostAlloc()`	分配页锁定主机内存	支持更快的主机-设备传输
`cudaMallocManaged()`	统一内存（托管内存）	由系统自动迁移数据，简化编程

常见误区：初学者常忽略 cudaMemcpy 的返回值。该函数返回 cudaError_t 类型，必须检查是否等于 cudaSuccess。否则，当显存不足或指针非法时，程序会静默失败，后续计算产生随机结果。建议封装错误检查宏：

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

4.5 案例复盘：性能瓶颈分析与优化

某数据中心运行一个矩阵乘法任务，初始实现直接使用三重循环核函数，性能远低于预期。通过 ncu（NVIDIA Nsight Compute）分析发现：

全局内存访问未合并：线程访问的地址不连续，导致每次内存事务仅利用 32 字节中的 4 字节，有效带宽仅为理论值的 12%。
缺乏共享内存复用：每个线程重复从全局内存读取矩阵元素，未利用线程块内的数据局部性。

优化策略：使用分块矩阵乘法（Tiling），将子矩阵加载到共享内存，并确保线程访问连续地址以满足内存合并访问条件。优化后有效带宽提升至理论值的 78%，整体性能提升 6.2 倍。

  复盘要点：
  先确保功能正确，再追求性能。使用 compute-sanitizer 检查越界与竞争条件。
内存拷贝是隐藏成本。尽量在设备端完成数据预处理，减少 H2D/D2H 传输次数。
核函数启动是异步的，但 cudaMemcpy 是同步的。使用 cudaStream 实现计算与拷贝重叠。
线程块大小通常设为 128~512 的倍数，需能被 32（warp 大小）整除。

CUDA 编程入门的关键在于建立“数据并行”思维：将问题分解为大量可独立执行的小任务，并合理组织线程层级与内存访问模式。掌握环境搭建、核函数编写和内存管理这三块基石后，即可进一步探索共享内存调优、流并发、Tensor Core 等高级主题。