第二章：CUDA编程基础——GPU架构与线程模型，CUDA C/C++基础语法，内存管理与数据传输，核函数编写与优化

好，我们直接进入正题。这一章是CUDA编程的基石，也是你后续做框架迁移时绕不开的核心。说白了，不懂GPU架构和线程模型，你写的核函数可能连CPU都不如。我见过太多人一上来就写并行代码，结果性能惨不忍睹——嗯，咱们先把地基打牢。

2.1 GPU架构与线程模型：从硬件到软件的逻辑映射

先聊聊GPU的硬件架构。你想想看，CPU是“大核少核”，适合处理复杂串行任务；GPU是“小核多核”，适合做数据并行计算。NVIDIA的GPU核心叫SM（Streaming Multiprocessor），每个SM里包含多个CUDA Core、共享内存、寄存器等资源。

我个人习惯把SM想象成一个“小工厂”，CUDA Core就是工厂里的工人。工人越多，并行能力越强。但要注意，工人不能独立干活，得按小组（Warp）来调度。一个Warp是32个线程，这是硬件调度的最小单位。

为什么会这样设计？因为GPU需要把大量线程映射到有限的SM上。Block里的线程可以协作，Grid里的Block则相对独立。我在项目中遇到过一个问题：Block开太大，导致共享内存不够用；开太小，又浪费了SM的计算资源。这个平衡点需要反复调。

2.2 CUDA C/C++基础语法：从Hello World到并行计算

CUDA C/C++其实就是在标准C/C++上加了一些扩展。你不需要学一门新语言，只需要记住几个关键字和API。

先看一个最简单的核函数：

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

这里有几个关键点：

• __global__：声明这是一个核函数，在设备上执行，从主机端调用
• threadIdx.x：当前线程在Block内的索引
• blockIdx.x：当前Block在Grid内的索引
• blockDim.x：每个Block包含的线程数

调用方式也很直接：

int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

注意那个尖括号语法，里面是Grid和Block的维度配置。我刚开始写的时候老搞混顺序——第一个是Grid维度，第二个是Block维度。嗯，这个记牢了。

2.3 内存管理与数据传输：别让数据搬运成为瓶颈

GPU编程里最容易被忽视的就是内存管理。你想想看，数据从CPU传到GPU，再从GPU传回来，这个传输过程走的是PCIe总线，带宽远低于显存带宽。我见过一个项目，90%的时间都花在数据传输上，真正的计算只占10%。

CUDA的内存模型分几种：

内存类型	位置	访问速度	作用域
全局内存	显存	慢（约200GB/s）	所有线程
共享内存	SM内部	快（约1TB/s）	Block内线程
寄存器	SM内部	极快	单个线程
常量内存	显存（带缓存）	快（只读）	所有线程

内存分配和传输的API很简单：

float *d_A;
cudaMalloc(&d_A, N * sizeof(float));  // 在GPU上分配内存
cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice);  // 从CPU拷贝到GPU
// ... 执行核函数 ...
cudaMemcpy(h_C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);  // 从GPU拷贝回CPU
cudaFree(d_A);  // 释放GPU内存

2.4 核函数编写与优化：从能跑到跑得快

写一个能跑的核函数不难，难的是让它跑得快。这里分享几个我常用的优化技巧。

技巧一：合并内存访问

GPU访问全局内存时，会以128字节为单位进行合并。如果相邻线程访问相邻地址，就能一次搞定。否则，就会产生多次内存事务，性能暴跌。

// 好的写法：线程0访问A[0]，线程1访问A[1]...
int i = threadIdx.x + blockIdx.x * blockDim.x;
C[i] = A[i] + B[i];

// 坏的写法：线程0访问A[0]，线程1访问A[1000]...
int i = threadIdx.x * 1000 + blockIdx.x;
C[i] = A[i] + B[i];

技巧二：使用共享内存减少全局内存访问

共享内存是SM内部的SRAM，速度比全局内存快一个数量级。如果Block内的线程需要反复访问同一块数据，把它搬到共享内存里。

__global__ void sharedMemExample(float* A, float* B, float* C, int N) {
    __shared__ float s_data[256];  // 共享内存数组
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    
    // 从全局内存加载到共享内存
    s_data[threadIdx.x] = A[i];
    __syncthreads();  // 同步，确保所有线程都加载完毕
    
    // 在共享内存上计算
    s_data[threadIdx.x] = s_data[threadIdx.x] * 2.0f;
    __syncthreads();
    
    // 写回全局内存
    C[i] = s_data[threadIdx.x];
}

技巧三：减少分支发散

一个Warp里的32个线程执行同一条指令。如果遇到if-else，部分线程走真分支，部分走假分支，Warp就会串行执行两个分支，性能减半。

// 避免这样写
if (i % 2 == 0) {
    C[i] = A[i] + B[i];
} else {
    C[i] = A[i] - B[i];
}

// 可以改成这样
C[i] = (i % 2 == 0) ? (A[i] + B[i]) : (A[i] - B[i]);  // 还是发散
// 更好的做法：重新组织数据，让同一个Warp里的线程走同一个分支

嗯，这里要注意：不是所有分支发散都能避免。有时候业务逻辑就是需要条件判断。那就尽量让同一个Warp里的线程条件一致，比如按奇偶索引来分组。

2.5 实战：向量加法完整示例

把上面学的东西串起来，写一个完整的向量加法程序：

#include <cuda_runtime.h>
#include <stdio.h>

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

int main() {
    int N = 1 << 20;  // 约100万个元素
    size_t size = N * sizeof(float);
    
    // 分配主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);
    
    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_A[i] = i * 1.0f;
        h_B[i] = i * 2.0f;
    }
    
    // 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);
    
    // 传输数据到GPU
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    
    // 配置并启动核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    
    // 等待GPU完成
    cudaDeviceSynchronize();
    
    // 传回结果
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
    
    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %f\n", i, h_C[i]);
    }
    
    // 清理
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(h_C);
    
    return 0;
}

这个例子虽然简单，但包含了CUDA编程的所有核心步骤：内存分配、数据传输、核函数配置、同步、结果回传。你在做框架迁移时，每一步都可能成为性能瓶颈。我个人习惯是先跑通功能，再用NVIDIA的Nsight工具做性能分析，看看瓶颈在哪。

下一章我们会深入CUDA流和事件，以及如何用它们实现异步计算和数据传输重叠。嗯，到时候你会看到，真正的性能优化才刚刚开始。