第2章算力引擎

算力引擎：GPU架构与通用计算（CUDA核心与Tensor Core详解）

大家好，我是老张。今天咱们聊聊GPU。

很多人觉得GPU就是显卡，打游戏用的。其实不然。在数据中心里，GPU早就是算力引擎的核心了。我2015年刚接触GPU通用计算时，还觉得这东西就是个“大号向量机”。后来做项目多了，才真正理解它的设计哲学。

说白了，GPU和CPU最大的区别是什么？CPU是“全能选手”，什么活都能干，但一次只能干几件。GPU是“流水线工人”，成千上万个线程一起上，专治各种并行计算。

1. CUDA核心：GPU的“小工”

CUDA核心是什么？你可以把它想象成一个非常简单的计算单元。它只能做加减乘除，没有复杂的分支预测，没有乱序执行。但好处是——它小啊！

一个GPU里有多少CUDA核心？拿NVIDIA的A100来说，有6912个。H100更是达到了18432个。你想想看，CPU才几个核？8个、16个顶天了。

但这里有个坑：CUDA核心不是独立工作的。它们被组织成SM（Streaming Multiprocessor，流式多处理器）。每个SM里有几十到上百个CUDA核心，共享指令缓存、共享内存和寄存器文件。

关键点：CUDA核心执行的是“单指令多线程”（SIMT）模型。也就是说，同一时刻，一个SM里的所有CUDA核心执行同一条指令，但处理不同的数据。

我在项目中遇到过一个问题：某个算法里有很多if-else分支。结果发现，GPU利用率只有30%。为什么？因为分支导致线程束（warp）里的线程走了不同路径，SM只能串行执行。这就是所谓的“线程束发散”。

避坑指南：我曾经因为没注意分支发散，导致一个图像处理项目性能惨不忍睹。后来把分支条件改成数据驱动的查找表，性能直接翻了3倍。记住：GPU讨厌分支，喜欢无脑计算。

2. Tensor Core：AI计算的“特种兵”

CUDA核心虽然多，但做矩阵乘法还是不够快。为什么？因为矩阵乘法需要大量乘加操作（FMA）。CUDA核心一次只能做一个FMA，效率太低。

于是NVIDIA在Volta架构里引入了Tensor Core。这东西是专门为深度学习设计的。它一次能做一个4x4矩阵乘法，也就是16个FMA操作。你想想看，效率提升了多少？

到了H100，Tensor Core已经进化到第四代了。支持FP8、FP16、BF16、TF32等多种精度。我最常用的是TF32——它用19位精度做计算，但输出结果可以累加到32位。说白了，就是“用更少的位宽，做更快的计算，精度损失还看不出来”。

架构	Tensor Core代数	支持精度	峰值TFLOPS（FP16）
Volta (V100)	第一代	FP16	125
Turing (T4)	第二代	FP16, INT8, INT4	65
Ampere (A100)	第三代	TF32, FP16, BF16, INT8, INT4	312
Hopper (H100)	第四代	FP8, FP16, BF16, TF32, INT8	1979

嗯，这里要注意：Tensor Core不是万能的。它只适合矩阵乘法、卷积这类规则计算。如果你的算法不规则，比如稀疏矩阵、图计算，Tensor Core就帮不上忙了。这时候还得靠CUDA核心。

3. 通用计算：CUDA编程模型

说了这么多硬件，咱们聊聊怎么用。CUDA编程模型的核心是三个词：线程、块、网格。

线程（Thread）：最小的执行单元，每个线程执行一个kernel函数实例。
线程块（Block）：一组线程，可以共享共享内存，可以同步。
网格（Grid）：一组线程块，构成整个计算任务。

我刚开始学CUDA时，总觉得这玩意儿抽象。后来想明白了：你把数据切成小块，每个块交给一个线程块处理。线程块里的线程协作处理这块数据。就这么简单。

来看个例子。假设我们要做向量加法：

__global__ void vecAdd(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() {
    // ... 分配内存、拷贝数据 ...
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
    // ... 拷贝结果回CPU ...
}

这段代码里，每个线程处理一个元素。blockIdx.x是块索引，threadIdx.x是线程在块内的索引。两者组合起来，就是全局线程ID。

个人经验：blockSize一般选256或512。太小了SM利用率低，太大了寄存器不够用。我习惯用256，因为它是32的倍数（一个warp是32线程），而且能整除大多数数据量。

4. 内存层次：别让数据成为瓶颈

GPU的性能瓶颈往往不在计算，而在内存。你想想看，CUDA核心算得再快，数据喂不进去也是白搭。

GPU的内存层次是这样的：

全局内存（Global Memory）：最大，但最慢。延迟400-800个周期。
共享内存（Shared Memory）：每个SM独有，很小但很快。延迟20-30个周期。
寄存器（Register）：每个线程独有，最快。延迟1个周期。
常量内存（Constant Memory）：只读，有缓存，适合存放常量。
纹理内存（Texture Memory）：有特殊缓存，适合空间局部性强的访问。

我记得有个项目是做矩阵转置。直接写全局内存，带宽利用率只有20%。后来用了共享内存做分块转置，利用率直接飙到80%。

为什么会这样？因为全局内存访问有“合并”要求。如果同一个warp里的线程访问连续地址，硬件会把多次访问合并成一次。否则，就是多次独立访问，带宽浪费严重。

核心原则：尽量让同一个warp里的线程访问连续的内存地址。这就是所谓的“内存合并访问”。

5. 实战：矩阵乘法优化

矩阵乘法是GPU的看家本领。咱们来看看怎么优化。

最简单的实现：每个线程算一个元素。但这样全局内存访问次数太多，性能很差。

优化思路：用共享内存做分块。每个线程块负责计算一个子矩阵。先把子矩阵的数据从全局内存加载到共享内存，然后从共享内存读取计算。

__global__ void matMulShared(float* A, float* B, float* C, int N) {
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
    
    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    for (int k = 0; k < N / BLOCK_SIZE; k++) {
        As[threadIdx.y][threadIdx.x] = A[row * N + k * BLOCK_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(k * BLOCK_SIZE + threadIdx.y) * N + col];
        __syncthreads();
        
        for (int i = 0; i < BLOCK_SIZE; i++) {
            sum += As[threadIdx.y][i] * Bs[i][threadIdx.x];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

这段代码里，__syncthreads()是同步点。确保所有线程都加载完数据后，再开始计算。否则，有的线程还在读，有的已经开始算了，数据就乱了。

注意：__syncthreads()只能在同一个线程块内同步。不同线程块之间无法同步。如果你需要全局同步，只能结束当前kernel，再启动下一个。

我做过测试：用共享内存优化后，矩阵乘法的性能提升了5-10倍。具体取决于矩阵大小和硬件配置。

6. Tensor Core编程：cuBLAS与cuDNN

直接编程Tensor Core很麻烦。NVIDIA提供了cuBLAS和cuDNN库，封装了底层细节。

比如用cuBLAS做矩阵乘法：

cublasHandle_t handle;
cublasCreate(&handle);

float alpha = 1.0f, beta = 0.0f;
cublasGemmEx(handle, CUBLAS_OP_N, CUBLAS_OP_N,
             M, N, K, &alpha,
             A, CUDA_R_16F, lda,
             B, CUDA_R_16F, ldb, &beta,
             C, CUDA_R_16F, ldc,
             CUBLAS_COMPUTE_32F,
             CUBLAS_GEMM_DEFAULT_TENSOR_OP);

cublasDestroy(handle);

这段代码里，cublasGemmEx会自动选择Tensor Core还是CUDA核心。你只需要指定精度和计算类型。CUBLAS_COMPUTE_32F表示用32位浮点累加，即使输入是FP16。

我个人习惯：能用库就用库。除非你有非常特殊的计算需求，否则不要自己写Tensor Core代码。库的优化比你手写的好得多。

7. 性能调优：从理论到实践

最后，分享几个性能调优的经验：

计算密度：确保每个线程有足够的计算量。如果计算太少，内存访问开销占比太大。
占用率（Occupancy）：每个SM上活跃的warp数量。占用率越高，越能隐藏内存延迟。
寄存器压力：每个线程用的寄存器太多，会降低占用率。我一般控制在32个以内。
内存带宽：用nvidia-smi或Nsight Systems监控带宽利用率。如果低于50%，说明内存访问模式有问题。

避坑指南：我曾经为了追求高占用率，把每个线程的寄存器压到16个。结果因为寄存器溢出（spill），性能反而下降了。后来发现，适当增加寄存器，减少全局内存访问，效果更好。记住：占用率不是越高越好，要综合考虑。

好了，关于GPU架构和通用计算，今天就聊到这儿。CUDA核心和Tensor Core各有千秋，关键是要根据你的应用场景选择合适的工具。做AI训练，Tensor Core是首选；做科学计算，CUDA核心更灵活。两者配合，才能发挥GPU的最大潜力。

公众号:蓝海资料掘金营，微信deep3321

📚 课程目录

算力引擎：GPU架构与通用计算（CUDA核心与Tensor Core详解）

1. CUDA核心：GPU的“小工”

2. Tensor Core：AI计算的“特种兵”

3. 通用计算：CUDA编程模型

4. 内存层次：别让数据成为瓶颈

5. 实战：矩阵乘法优化

6. Tensor Core编程：cuBLAS与cuDNN

7. 性能调优：从理论到实践