GPU 并行计算原理

GPU 并行计算的核心在于其大规模并行处理能力。本文将深入探讨 GPU 并行计算的原理、模型和优化策略。

SIMT（单指令多线程）模型

GPU 采用 SIMT（Single Instruction, Multiple Threads）架构，这是一种特殊的并行计算模型：

SIMT 的工作原理

单指令：Warp 中的所有线程同时执行相同的指令
多线程：每个线程处理不同的数据
数据并行：通过数据并行实现高性能计算

// SIMT 示例：向量加法
__global__ void vectorAdd(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];  // 所有线程执行相同操作
    }
}

线程分歧（Warp Divergence）

当 Warp 内的线程需要执行不同的代码路径时，会产生线程分歧：

// 线程分歧示例
__global__ void divergentKernel(int *data) {
    int idx = threadIdx.x;
    if (idx % 2 == 0) {
        data[idx] = idx * 2;  // 偶数线程
    } else {
        data[idx] = idx * 3;  // 奇数线程
    }
}

线程分歧会导致性能下降，因为 GPU 需要串行执行不同的分支。优化方法包括：

尽量减少条件分支
使用 warp 内建函数（如 __any、__all）
重新组织数据使相同条件的线程相邻

内存合并访问

内存合并访问是 GPU 性能优化的关键概念：

什么是内存合并

当相邻线程访问连续的内存地址时，GPU 可以将多个内存请求合并为一个事务，提高带宽利用率。

// 好的内存合并
__global__ void coalescedAccess(float *data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] = idx;  // 连续访问
}

// 不好的内存访问
__global__ void stridedAccess(float *data, int stride) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx * stride] = idx;  // 跨步访问
}

内存合并的最佳实践

使用行优先存储（Row-major）
避免跨步访问（Strided Access）
使用共享内存进行数据重排
考虑数据对齐（Alignment）

共享内存优化

共享内存是 GPU 上最快的可编程存储器，合理使用可以显著提升性能。

共享内存的特性

速度快：比全局内存快 100 倍以上
容量小：通常 48KB-96KB
Block 共享：同一 Block 内的线程共享
可编程：完全由程序员控制

矩阵乘法示例

#define TILE_SIZE 16

__global__ void matrixMul(float *A, float *B, float *C, int N) {
    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
    __shared__ float tileB[TILE_SIZE][TILE_SIZE];

    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;

    int row = by * TILE_SIZE + ty;
    int col = bx * TILE_SIZE + tx;

    float sum = 0.0f;

    for (int t = 0; t < N / TILE_SIZE; ++t) {
        tileA[ty][tx] = A[row * N + t * TILE_SIZE + tx];
        tileB[ty][tx] = B[(t * TILE_SIZE + ty) * N + col];
        __syncthreads();  // 等待所有线程完成加载

        for (int k = 0; k < TILE_SIZE; ++k) {
            sum += tileA[ty][k] * tileB[k][tx];
        }
        __syncthreads();  // 等待所有线程完成计算
    }

    C[row * N + col] = sum;
}

占用率（Occupancy）

占用率是指 SM 上活跃 Warp 的数量与最大 Warp 数量的比值。

影响占用率的因素

寄存器使用：每个线程使用的寄存器数量
共享内存使用：每个 Block 使用的共享内存
Block 大小：每个 Block 的线程数

优化占用率

// 检查占用率的 CUDA 工具
// nvcc --ptxas-options=-v kernel.cu

// 减少寄存器使用
__global__ void __launch_bounds__(256, 2) optimizedKernel(int *data) {
    // 使用较少的寄存器
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] = idx;
}

原子操作

当多个线程需要同时访问同一内存位置时，需要使用原子操作。

// 原子操作示例
__global__ void histogram(int *data, int *hist, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        int val = data[idx];
        atomicAdd(&hist[val], 1);  // 原子加法
    }
}

// 其他原子操作
// atomicSub, atomicExch, atomicMin, atomicMax
// atomicInc, atomicDec, atomicCAS, atomicAnd, atomicOr, atomicXor

流（Streams）

流允许在 GPU 上并发执行多个操作，提高硬件利用率。

// 使用流实现计算和传输重叠
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

float *d_data1, *d_data2;
cudaMalloc(&d_data1, N * sizeof(float));
cudaMalloc(&d_data2, N * sizeof(float));

// 在不同流中异步执行
kernel<<>>(d_data1);
kernel<<>>(d_data2);

// 等待所有流完成
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

性能分析工具

Nsight Compute：详细的 GPU 性能分析
Nsight Systems：系统级性能分析
nvprof：命令行性能分析工具
Visual Profiler：图形化性能分析

总结

GPU 并行计算通过 SIMT 模型、内存合并、共享内存优化等技术实现高性能。掌握这些原理和优化技巧，可以充分发挥 GPU 的计算能力，为科学计算、AI 训练等应用提供强大的支持。