CUDA 编程基础回顾

好，咱们正式开始。这个系列课程讲的是多GPU负载均衡，但上来我得先带你把单卡上的CUDA基础再过一遍。别嫌啰嗦——我在实际调优时发现，大部分多卡性能瓶颈其实都根植于单卡基础没打牢。

1. GPU架构概览

先说说GPU长什么样。你想想看，CPU是“精壮小子”：核少但单个核又大又强，擅长复杂控制流。GPU呢？它是个“人海战术”的典范：成百上千个轻量级核心堆在一起，专攻数据并行。

我记得早期用Fermi架构时，一个SM只有32个CUDA核心，到了现在的H100，一个SM可以塞进128个。但不管怎么变，Warp始终是调度的基本单位——32个线程一捆，一起执行同一指令。这是理解性能的关键。

架构代号	SM特性	典型算力	我的踩坑点
Fermi	16个SM，每个32 CUDA核	1.3 TFLOPS	共享内存没bank冲突时很开心，有冲突就血崩
Kepler	SMX，动态并行	4 TFLOPS	动态并行听着牛，实际调试起来想哭
Volta	独立线程调度，Tensor Core	7.5 TFLOPS	线程束发散第一次把我整懵了
Ampere/Hopper	MIG，DPX指令	9~60 TFLOPS	多实例GPU切分时资源隔离要小心

2. CUDA核心概念

说白了，CUDA编程就是告诉GPU：你的数据在哪儿、怎么切、每个小片用哪个函数处理。这里有两个主角：Host（CPU端）和Device（GPU端）。

• Kernel：你写的一个 __global__ void 函数，在设备上并行执行。
• 流（Stream）：一系列按顺序执行的命令队列。多流可以并发。
• 事件（Event）：用于计时或同步，比 cudaDeviceSynchronize 更精细。

我习惯在写Kernel前先画一个数据流水线图——从Host搬运到Device，算完再搬回来。这一步看似简单，但曾经有个项目我把两批独立数据放进了同一个流，结果吞吐量直接腰斩。

3. 线程层次结构

CUDA的线程组织是三层嵌套：Grid → Block → Thread。每个Block内的线程可以协作（通过共享内存和同步原语），不同Block之间则是独立的。

// 典型的Kernel启动配置
dim3 gridDim(256, 1, 1);   // 网格中有256个block
dim3 blockDim(128, 1, 1);  // 每个block中有128个thread

myKernel<<<gridDim, blockDim>>>(d_data, N);

Block内的线程数量最好是Warp大小（32）的整数倍，不然会有空着的线程束浪费资源。嗯，这里要注意：一个Block最多1024个线程（计算能力3.0+），但别真塞满——要留给寄存器、共享内存总量。

4. 内存层次

CUDA的内存模型，我把它总结成一张“速度‐容量”图。说白了：离计算单元越近越快，但越小越贵。

内存类型	作用域	典型容量	延迟（时钟周期）	缓存？
寄存器（Register）	单个线程	每个线程最多255个	~1	无
局部内存（Local Memory）	单个线程	依赖寄存器溢出	~400	有L1/L2
共享内存（Shared Memory）	整个Block	每个SM通常48KB～164KB	~5（无bank冲突）	无，手动管理
全局内存（Global Memory）	所有线程 + Host	卡上显存（GB级）	~400	有L2
常量内存（Constant Memory）	所有线程，只读	64KB	~5（缓存命中时）	专用常量缓存
纹理内存（Texture Memory）	所有线程，只读	同显存	~400（无缓存）	专用纹理缓存

我刚开始做GPU加速时，一股脑把数据全扔进全局内存，然后被慢得怀疑人生。后来才发现，很多场景下把频繁访问的数据搬进共享内存，性能能翻几番。

5. Kernel编写实战

光说不练假把式。咱们写一个最经典的向量加法Kernel，顺便把上面讲的概念串起来。

// vector_add.cu
#include <cuda_runtime.h>

__global__ void vecAddKernel(const float* A, const float* B, float* C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

int main() {
    int N = 1 << 20;  // 约100万
    size_t bytes = N * sizeof(float);

    float *h_A, *h_B, *h_C;
    float *d_A, *d_B, *d_C;

    // 1. 分配Host内存
    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. 分配Device内存
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);

    // 3. 数据从Host搬到Device
    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;
    vecAddKernel<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

    // 5. 结果从Device搬回Host
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

    // 6. 验证
    for (int i = 0; i < 10; i++) {
        printf("h_C[%d] = %f\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;
}

你看，这个流程其实就六步：分配Host内存 → 分配Device内存 → 数据传输 → 启动Kernel → 传回结果 → 清理。简单吧？但我在实际项目中踩过坑：忘了检查cudaMalloc的返回值，结果在大卡上分配了过大的共享内存直接挂掉。所以——

好了，基础回顾到这儿。下面这张SVG图把这一章的核心脉络画出来了，你留个印象。