第一章：CUDA入门——从GPU架构到第一个程序跑起来

大家好，我是你们这门课的主讲。做了十几年GPU计算，踩过不少坑，也积累了一些经验。今天咱们聊聊CUDA入门最核心的几个东西：GPU到底长什么样、CUDA怎么跟GPU对话、以及怎么把开发环境搭好。

说实话，很多初学者一上来就写kernel，结果性能惨不忍睹。为什么？因为你不了解GPU的脾气。好，咱们从头捋一遍。

1.1 GPU架构概述——它为什么这么快？

GPU和CPU的设计哲学完全不同。CPU是「全能选手」，什么活儿都能干，但一次干不了太多。GPU是「流水线工人」，成千上万个核心一起上，专攻并行计算。

我当年第一次看NVIDIA的架构白皮书时，最震撼的是它的层次结构。咱们以目前主流的Ampere架构（比如A100）为例，简单拆解一下：

• SM（Streaming Multiprocessor）：这是GPU的核心计算单元。一个GPU有几十到上百个SM。
• CUDA Core：每个SM里有几十到上百个CUDA核心。这些才是真正干活的「小兵」。
• 共享内存（Shared Memory）：SM内部的高速缓存，延迟极低。我项目中很多优化都是靠它。
• 全局内存（Global Memory）：显存，容量大但慢。访问一次要几百个时钟周期。
• 寄存器（Registers）：每个线程私有的最快存储，但数量有限。

嗯，这里要注意：不是线程越多越好。每个线程能用的寄存器、共享内存都是有限的。我曾经有个项目，线程数翻了一倍，性能反而掉了一半——因为寄存器不够用，数据被挤到全局内存里去了。

1.2 CUDA编程模型——你怎么指挥GPU干活？

说白了，CUDA编程就是「你告诉GPU：这里有N个数据，你派M个线程去处理」。但怎么组织这M个线程，是有讲究的。

CUDA把线程组织成三层结构：

• 线程（Thread）：最小的执行单元。每个线程执行相同的kernel函数，但处理不同的数据。
• 线程块（Block）：一组线程，可以协作。同一个block内的线程可以通过共享内存通信。
• 网格（Grid）：一组block，构成整个任务。

你想想看，这就像公司组织架构：Grid是公司，Block是部门，Thread是员工。部门内部（Block内）可以开小会（共享内存），跨部门就得走邮件（全局内存），慢很多。

来看一个最简单的CUDA程序——向量加法。这玩意儿虽然简单，但麻雀虽小五脏俱全：

// 文件名：vector_add.cu
#include <stdio.h>

// GPU kernel：每个线程处理一个元素
__global__ void vecAdd(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1 << 20;  // 1048576个元素
    size_t bytes = n * sizeof(float);

    // 1. 分配主机内存
    float *h_a = (float*)malloc(bytes);
    float *h_b = (float*)malloc(bytes);
    float *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. 把数据从CPU拷贝到GPU
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 4. 启动kernel
    int blockSize = 256;
    int gridSize = (n + blockSize - 1) / blockSize;
    vecAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

    // 5. 把结果拷回CPU
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // 验证结果
    printf("c[0] = %f\n", h_c[0]);  // 应该输出 3.0

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

    return 0;
}

这段代码里，有几点我想强调一下：

1. __global__ 关键字告诉编译器，这个函数跑在GPU上，从CPU调用。
2. threadIdx、blockIdx、blockDim 是CUDA内置变量，帮你算当前线程的全局ID。
3. cudaMalloc和cudaMemcpy 是显式内存管理——GPU不会自动帮你搬数据。

1.3 CUDA开发环境搭建——磨刀不误砍柴工

好，理论说完了，咱们动手搭环境。我建议按这个顺序来：

第一步：安装NVIDIA驱动

驱动是GPU和操作系统之间的桥梁。去NVIDIA官网下载对应你显卡型号的最新驱动。装完后在命令行敲 nvidia-smi，如果能看到显卡信息，就说明驱动装好了。

我个人习惯用Linux（Ubuntu）做开发，Windows也行，但有些profiling工具在Linux下更顺手。

第二步：安装CUDA Toolkit

CUDA Toolkit包含了编译器（nvcc）、库（如cuBLAS、cuFFT）和调试工具。我建议用runfile安装，这样你能控制安装路径。装完后设置环境变量：

export PATH=/usr/local/cuda-12.0/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/cuda-12.0/lib64:$LD_LIBRARY_PATH

验证安装：nvcc --version，看到版本号就对了。

第三步：安装Nsight工具

Nsight是NVIDIA的性能分析神器。我强烈建议你从第一天就开始用它。Nsight Systems做宏观分析，Nsight Compute做微观kernel分析。

装好之后，跑一下刚才的向量加法，用Nsight看看：

nsys profile -o vec_add_report ./vector_add
ncu --target-processes all ./vector_add

你会看到kernel执行时间、内存带宽利用率、occupancy等指标。嗯，这些数字以后会是你优化路上的好朋友。

组件	版本建议	验证命令
NVIDIA驱动	最新稳定版（≥525）	nvidia-smi
CUDA Toolkit	12.x（与驱动兼容）	nvcc --version
Nsight Systems	2024.x	nsys --version
Nsight Compute	2024.x	ncu --version

小结

这一章咱们把GPU架构、CUDA编程模型和环境搭建都过了一遍。你想想看，其实核心就三件事：理解硬件层次、学会组织线程、搭好开发环境。后面的章节，我们会在这个基础上，一步步深入——从内存优化到算子融合，从profiling到性能调优。

好，第一章就到这里。记住：别急着写复杂代码，先把环境玩熟，把最简单的vector add跑通、跑明白。我曾经带过一个新人，上来就写卷积算子，结果环境都没配好，折腾了两周。嗯，咱们不犯这个错。

想要更多资料，联系可可