3、GPU加速实战：CUDA编程模型、Tensor Core原理、内存优化技巧

好，咱们进入GPU加速这个硬核话题。说实话，很多做AI加速的同学，一开始都卡在“怎么让GPU真正跑起来”这个坎上。你想想看，一张A100放在服务器里，跑出来的吞吐量还不如人家优化过的V100，这种事儿我见过太多次了。

这一章，我会把CUDA编程模型、Tensor Core原理、还有内存优化这三个核心点串起来讲。不是那种“Hello World”式的教学，而是真正能在实时推理场景里帮你提性能的实战经验。

3.1 CUDA编程模型——别只把它当“多线程”

很多人学CUDA，上来就背Grid、Block、Thread的层级关系。嗯，背这个没错，但你要理解它背后的设计意图。

CUDA的编程模型，说白了就是“分层并行”。

• Grid：对应整个任务。比如你要处理一张1024x1024的图片，这就是一个Grid。
• Block：对应任务的一个子块。每个Block内部，线程可以共享数据（Shared Memory）。
• Thread：最细粒度的执行单元。每个Thread处理一个像素或者一个元素。

我在项目中遇到过一个问题：有人把Block设得特别大，比如512x512，结果每个Block里的Shared Memory不够用，性能反而下降了。后来我改成16x16的Block，每个Block里256个线程，刚好把数据塞进Shared Memory，吞吐量直接翻倍。

来看一个简单的向量加法示例，感受一下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() {
    // ... 分配内存、拷贝数据 ...
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    // ... 拷贝结果回Host ...
}

你看，核心逻辑就三行。但真正要优化的，是后面的内存访问模式和数据搬运策略。

3.2 Tensor Core原理——AI加速的秘密武器

从Volta架构开始，NVIDIA引入了Tensor Core。这东西不是普通的CUDA Core，它是专门为矩阵乘法设计的硬件单元。

为什么Tensor Core快？因为它能在一个时钟周期内完成一个4x4矩阵的乘加运算（D = A * B + C）。而普通CUDA Core需要多个周期才能完成同样的事。

我建议你记住这个表格：

特性	CUDA Core	Tensor Core
计算单元	标量	4x4矩阵
精度支持	FP32/FP64/INT32	FP16/BF16/INT8/INT4
吞吐量（A100）	19.5 TFLOPS (FP32)	312 TFLOPS (FP16)
适用场景	通用计算	矩阵乘法、卷积

看到没？Tensor Core的吞吐量是CUDA Core的十几倍。但有个前提——你的数据精度得降到FP16或INT8。

我曾经在一个实时语音识别项目里，把模型从FP32量化到FP16，然后用Tensor Core做矩阵乘法。推理延迟从12ms降到了3ms，而且精度损失几乎可以忽略。嗯，这里要注意：量化不是无脑降精度，你得做校准，不然某些敏感层会崩。

3.3 内存优化技巧——别让数据搬运成为瓶颈

GPU计算再快，如果数据搬来搬去，性能一样拉胯。我见过最夸张的情况：一个模型在GPU上计算只花了2ms，但数据从CPU拷贝到GPU花了15ms。你说这加速了个寂寞？

内存优化的核心就三个字：少搬运。

具体怎么做？我总结了几个实战经验：

• 使用Pinned Memory（页锁定内存）：用 cudaHostAlloc 分配Host内存，这样CPU到GPU的拷贝带宽能提升2~3倍。
• 异步传输：用 cudaMemcpyAsync 配合Stream，让计算和拷贝重叠。比如你在GPU算第N批数据的同时，把第N+1批数据拷进来。
• Shared Memory优化：把频繁访问的数据放到Shared Memory里，避免反复访问Global Memory。Global Memory的延迟是400~800个周期，Shared Memory只有几十个周期。
• 合并访问（Coalesced Access）：让相邻的线程访问相邻的内存地址。这个在矩阵转置、卷积操作里特别重要。

3.4 实战：一个完整的CUDA优化流程

咱们来看一个具体的例子。假设你要在GPU上做一个矩阵乘法：C = A * B，其中A是1024x1024，B是1024x1024。

第一步，用Naive的CUDA实现：

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

这个实现，每个线程都要从Global Memory读N次A和N次B，总共2*N^3次访存。对于1024x1024的矩阵，就是20亿次访存。太慢了。

第二步，用Shared Memory优化：

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

这个版本，每个线程只从Global Memory读两次数据（一次A，一次B），然后从Shared Memory里反复读。访存量从2*N^3降到了2*N^2。对于1024x1024的矩阵，就是200万次访存。快了100倍。

第三步，如果数据精度允许，用Tensor Core：

// 使用cuBLAS的Tensor Core版本
cublasGemmEx(handle, CUBLAS_OP_N, CUBLAS_OP_N,
             N, N, N,
             &alpha, A, CUDA_R_16F, N,
             B, CUDA_R_16F, N,
             &beta, C, CUDA_R_16F, N,
             CUDA_R_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP);

这个版本，直接调用cuBLAS的Tensor Core接口，吞吐量能到312 TFLOPS。比Shared Memory版本又快了一个数量级。

你看，从50到3000，差了60倍。这就是优化的力量。

3.5 知识结构图

下面这张图，把CUDA编程模型、Tensor Core原理、内存优化技巧串在了一起。你可以把它当作本章的思维导图。

这张图里，三个分支是并列关系。实际项目中，你往往需要同时用上这三种技术。比如，先用CUDA编程模型设计好并行结构，再用Tensor Core做矩阵运算，最后用内存优化技巧减少数据搬运。三者缺一不可。