GPU架构与CUDA基础：从硬件到编程模型

核心原理：GPU为何适合并行计算

GPU（图形处理器）最初为图形渲染设计，其架构天然偏向大规模并行。与CPU追求单线程低延迟不同，GPU将大量晶体管用于计算单元而非控制逻辑与缓存。以NVIDIA A100为例，它包含108个流式多处理器（SM），每个SM内有64个CUDA核心，总计6912个FP32核心。这种“多核心、宽SIMD”的设计，使得GPU在吞吐量（throughput）上远超CPU，尤其适合矩阵乘法、卷积等数据并行任务。

理解GPU架构的关键在于三个层次：全局层面（GPU芯片→SM阵列）、SM内部（CUDA核心、共享内存、调度器）、线程执行模型（warp调度）。下图展示了从CUDA线程到GPU硬件的映射关系，这是理解后续并行编程的基础。

图中左侧是CUDA软件层次，右侧是GPU硬件资源。关键映射关系为：Grid→GPU设备，Block→SM，Warp（32线程）→SM内的调度单元。一个Block只会被调度到同一个SM上，因此Block内的线程可以通过共享内存协作；而不同Block之间无法直接通信，只能通过全局内存间接交互。

CUDA编程模型：线程层次与内存层次

CUDA（Compute Unified Device Architecture）是NVIDIA提供的并行计算平台。其编程模型围绕三个核心抽象展开：线程层次、内存层次和同步机制。一个典型的CUDA程序由主机端（CPU）代码和设备端（GPU）内核函数（kernel）组成。

线程层次结构

• Thread（线程）：最小的执行单元，每个线程有独立的寄存器文件和程序计数器。
• Warp（线程束）：硬件调度单位，由32个线程组成。SIMD方式执行——同一时刻执行同一条指令。
• Block（线程块）：一组线程（通常128~512个），可包含多维结构（如二维blockDim(16,16)）。Block内线程可通过__syncthreads()同步。
• Grid（网格）：内核启动时创建的所有Block集合，Grid内Block之间无同步机制。

内核启动时需指定Grid和Block维度，例如：

// 定义内核函数
__global__ void vectorAdd(float *A, float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

// 主机端启动内核
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

这里blockIdx.x、threadIdx.x是内置变量，分别表示当前Block在Grid中的索引和当前线程在Block中的索引。blockDim.x表示Block的维度大小。这种一维索引方式是最常见的模式，实际中也可使用二维或三维索引处理图像等数据。

内存层次详解

CUDA内存模型与硬件层次紧密对应，理解各内存的作用域、生命周期和访问速度是优化性能的关键。下表总结了主要内存类型的特点：

共享内存是CUDA性能优化的核心资源。它位于SM内部，容量有限（A100为192KB/SM），但访问速度远快于全局内存。典型优化策略是分块（tiling）：将数据分块加载到共享内存，Block内线程反复使用这些数据，减少全局内存访问次数。

方法流程：CUDA程序开发与执行步骤

一个完整的CUDA程序从编写到执行，通常遵循以下流程：

1. 环境准备：安装NVIDIA驱动、CUDA Toolkit，使用nvcc编译器。
2. 内存分配：使用cudaMalloc为设备端分配全局内存，cudaMemcpy进行主机与设备间的数据拷贝。
3. 内核编写与启动：定义__global__函数，通过<<<grid, block>>>语法启动。
4. 同步与错误检查：调用cudaDeviceSynchronize()等待内核完成，使用cudaGetLastError()捕获错误。
5. 结果回传与释放：将计算结果拷贝回主机，调用cudaFree释放设备内存。

以下是一个完整的向量加法示例，展示了上述流程：

#include <cuda_runtime.h>
#include <iostream>

__global__ void vectorAdd(float *A, 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;  // 1M元素
    size_t bytes = N * sizeof(float);
    
    // 1. 分配主机内存并初始化
    float *h_A = new float[N];
    float *h_B = new float[N];
    float *h_C = new float[N];
    for (int i = 0; i < N; i++) { h_A[i] = i * 0.5f; h_B[i] = i * 0.3f; }
    
    // 2. 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);
    
    // 3. 拷贝数据到设备
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
    
    // 4. 启动内核
    int threads = 256;
    int blocks = (N + threads - 1) / threads;
    vectorAdd<<<blocks, threads>>>(d_A, d_B, d_C, N);
    
    // 5. 同步并检查错误
    cudaDeviceSynchronize();
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        std::cerr << "CUDA error: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }
    
    // 6. 回传结果并验证
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
    for (int i = 0; i < 10; i++) {
        std::cout << h_C[i] << " ";
    }
    std::cout << std::endl;
    
    // 7. 释放资源
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    delete[] h_A; delete[] h_B; delete[] h_C;
    return 0;
}

案例场景：矩阵乘法优化中的内存层次应用

矩阵乘法（C = A × B）是深度学习中最核心的运算之一。以M=N=K=1024的方阵为例，朴素实现中每个线程计算C矩阵的一个元素，需要读取A的一行（1024个float）和B的一列（1024个float），共2048次全局内存访问。整个计算需要1024×1024×2048×4字节 ≈ 8GB的全局内存流量，而实际计算量仅为2×1024³ ≈ 2.1GFLOP，内存带宽成为瓶颈。

采用共享内存分块优化后，将矩阵划分为16×16的Tile：

__global__ void matMulShared(float *A, float *B, float *C, int N) {
    __shared__ float As[16][16];
    __shared__ float Bs[16][16];
    
    int row = blockIdx.y * 16 + threadIdx.y;
    int col = blockIdx.x * 16 + threadIdx.x;
    float sum = 0.0f;
    
    for (int tile = 0; tile < N/16; tile++) {
        // 协作加载Tile到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + tile * 16 + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(tile * 16 + threadIdx.y) * N + col];
        __syncthreads();  // 确保加载完成
        
        // 计算当前Tile的部分和
        for (int k = 0; k < 16; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();  // 防止覆盖共享内存
    }
    C[row * N + col] = sum;
}

优化后，每个Tile只需从全局内存读取16×16×2×4字节 = 2KB数据，而计算量为2×16³ = 8192次乘加。全局内存访问量降低到原来的1/16，计算密度大幅提升。实际测试中，这种优化在A100上可将矩阵乘法性能从约1.2 TFLOPS提升至19.5 TFLOPS，接近硬件峰值。

误区与复盘：常见CUDA编程陷阱

在实际开发中，以下问题频繁出现，值得特别警惕：

另一个常见误区是忽视内存对齐。全局内存访问时，如果每个线程访问的地址不满足128字节对齐（如使用float但偏移量不是4的倍数），会导致多次内存事务，降低有效带宽。使用cudaMallocPitch或cudaMalloc3D可以自动处理对齐问题。

复盘一个实际案例：某团队在训练Transformer模型时，发现GPU利用率仅35%。通过ncu（NVIDIA Nsight Compute）分析发现，主要瓶颈是cudaMemcpy频繁同步拷贝——每个step都从设备回传loss到主机。优化方案是使用cudaMemcpyAsync配合cudaStream实现异步拷贝，将数据传输与下一轮计算重叠，最终GPU利用率提升至82%。

最后需要强调的是，warp发散问题。当warp内32个线程执行不同分支时，所有路径会串行执行。例如：

if (idx % 2 == 0) {
    // 偶数路径
} else {
    // 奇数路径
}

这段代码在warp中会导致两个分支都执行，效率减半。优化策略是尽量让分支条件基于warp级别的粒度（如warpId % 2），或使用三元运算符替代分支。

掌握GPU架构与CUDA基础，是后续理解多GPU并行（如NCCL通信、模型并行）的基石。只有深入理解线程映射、内存层次和同步机制，才能写出高效、可扩展的分布式训练代码。