GPU编程模型

编程模型概述

GPU采用SIMT (Single Instruction Multiple Threads) 执行模型， 与CPU的SIMD模型有本质区别。理解这一差异是编写高效GPU代码的基础。

CPU vs GPU架构特点：
┌──────────────────────────────────────────┐
│              CPU           GPU           │
├──────────────────────────────────────────┤
│ 核心数       少(8-128)     多(数千)      │
│ 单核性能     高            低            │
│ 缓存大小     大(30-50MB)   小(4-6MB)     │
│ 内存带宽     低(50GB/s)    高(2TB/s+)    │
│ 延迟容忍度   低            高            │
│ 适用场景     复杂逻辑      并行计算      │
│ 编程模型     MIMD          SIMT          │
└──────────────────────────────────────────┘

SIMT执行模型：
├── Warp (线程束)
│   ├── 32个线程为一组
│   ├── 执行相同指令
│   └── 分支发散影响性能
│
├── 线程块 (Block)
│   ├── 多个Warp组成
│   ├── 共享内存可见
│   └── 可同步
│
└── 网格 (Grid)
    ├── 多个Block组成
    ├── Block间独立执行
    └── 全局内存可见

CUDA编程模型

Kernel函数

// CUDA Kernel函数示例
__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];
    }
}

// Kernel调用
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

// Kernel启动配置
// <<<gridDim, blockDim, sharedMemSize, stream>>>
// gridDim: 网格维度 (dim3)
// blockDim: 线程块维度 (dim3)
// sharedMemSize: 动态共享内存大小
// stream: CUDA流

线程层次

线程层次结构：
┌──────────────────────────────────────────┐
│ Grid (网格)                              │
│ ├── gridDim.x, gridDim.y, gridDim.z      │
│ ├── blockIdx.x, blockIdx.y, blockIdx.z   │
│ └── 包含多个Block                        │
│                                          │
│   └── Block (线程块)                     │
│       ├── blockDim.x, blockDim.y, .z     │
│       ├── threadIdx.x, threadIdx.y, .z   │
│       ├── 最大1024线程                   │
│       └── 共享内存可见                   │
│                                          │
│         └── Warp (线程束)                │
│             ├── 32个连续线程              │
│             └── 执行相同指令              │
└──────────────────────────────────────────┘

全局索引计算示例：
// 1D Grid, 1D Block
int idx = blockIdx.x * blockDim.x + threadIdx.x;

// 2D Grid, 2D Block
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = blockIdx.y * blockDim.y + threadIdx.y;
int index = idy * width + idx;

// 3D Grid, 3D Block
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = blockIdx.y * blockDim.y + threadIdx.y;
int idz = blockIdx.z * blockDim.z + threadIdx.z;
int index = idz * width * height + idy * width + idx;

内存模型

全局内存

全局内存特点：
├── 容量大 (40-80GB)
├── 延迟高 (~300-500 cycles)
├── 带宽高 (1-3TB/s)
└── 所有线程可见

内存合并访问：
├── 同一Warp线程访问连续内存
├── 对齐访问 (128B对齐)
└── 合并为少数内存事务

// 合并访问示例 (高效)
__global__ void copy(float *dst, float *src, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        dst[idx] = src[idx];  // 连续访问
    }
}

// 跨步访问 (低效)
__global__ void copyStride(float *dst, float *src, int n, int stride) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx * stride < n) {
        dst[idx] = src[idx * stride];  // 跨步访问
    }
}

// 随机访问 (最差)
__global__ void gather(float *dst, float *src, int *indices, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        dst[idx] = src[indices[idx]];  // 随机访问
    }
}

共享内存

共享内存特点：
├── 容量小 (48-228KB per SM)
├── 延迟低 (~20-30 cycles)
├── 带宽高 (~数TB/s)
└── Block内线程共享

// 矩阵乘法优化示例
__global__ void matmulShared(float *A, float *B, float *C, 
                              int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 共享内存缓存
    __shared__ float As[16][16];
    __shared__ float Bs[16][16];
    
    float sum = 0.0f;
    
    // 分块计算
    for (int t = 0; t < K; t += 16) {
        // 加载到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * K + t + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t + threadIdx.y) * N + col];
        __syncthreads();
        
        // 块内矩阵乘
        for (int k = 0; k < 16; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();
    }
    
    C[row * N + col] = sum;
}

共享内存Bank冲突：
├── 32个Bank
├── 连续4字节映射到连续Bank
├── 多线程访问同一Bank导致冲突
└── 使用Padding避免冲突

性能优化

GPU性能优化核心策略：
├── 最大化并行度
│   ├── 充分利用所有SM
│   ├── 足够的线程块数量
│   └── 避免线程束发散
│
├── 优化内存访问
│   ├── 合并全局内存访问
│   ├── 使用共享内存缓存
│   └── 避免Bank冲突
│
├── 隐藏延迟
│   ├── 增加线程级并行
│   ├── 使用多流
│   └── 指令级并行
│
└── 减少数据传输
    ├── 减少Host-Device传输
    ├── 使用Pinned Memory
    └── 异步传输

常见优化技巧：
├── 循环展开 (#pragma unroll)
├── 使用快速数学函数 (__sinf, __expf)
├── 避免原子操作竞争
├── 使用常量内存
└── 使用纹理内存

Tensor Core

Tensor Core简介：
├── 专用矩阵计算单元
├── 单周期完成矩阵乘累加
├── 支持混合精度 (FP16/BF16/TF32/FP8)
└── 大幅提升AI计算性能

// 使用WMMA API (CUDA C++)
#include <mma.h>
using namespace nvcuda::wmma;

__global__ void tensorCoreGemm(half *A, half *B, float *C,
                                int M, int N, int K) {
    // 定义片段
    fragment<matrix_a, 16, 16, 16, half, row_major> a_frag;
    fragment<matrix_b, 16, 16, 16, half, col_major> b_frag;
    fragment<accumulator, 16, 16, 16, float> c_frag;
    
    // 初始化累加器
    fill_fragment(c_frag, 0.0f);
    
    // 加载矩阵块
    load_matrix_sync(a_frag, A, K);
    load_matrix_sync(b_frag, B, K);
    
    // Tensor Core矩阵乘
    mma_sync(c_frag, a_frag, b_frag, c_frag);
    
    // 存储结果
    store_matrix_sync(C, c_frag, N, mem_row_major);
}

// 使用cuBLAS调用Tensor Core
cublasGemmEx(handle, 
    CUBLAS_OP_N, CUBLAS_OP_N,
    M, N, K,
    &alpha, A, CUDA_R_16F, K,
            B, CUDA_R_16F, N,
    &beta,  C, CUDA_R_32F, N,
    CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP);

高级主题

高级GPU编程技术：
├── CUDA流 (Stream)
│   ├── 异步执行
│   ├── 多流并行
│   └── 事件同步
│
├── 动态并行
│   ├── Kernel内启动Kernel
│   └── 递归算法支持
│
├── Cooperative Groups
│   ├── 灵活的线程协作
│   ├── 跨Block同步
│   └── Grid级操作
│
├── CUDA Graph
│   ├── 图执行模型
│   ├── 减少启动开销
│   └── 适合重复工作流
│
└── Multi-Process Service (MPS)
    ├── 多进程GPU共享
    ├── 提高GPU利用率
    └── 适合推理服务

// 多流示例
cudaStream_t streams[4];
for (int i = 0; i < 4; i++) {
    cudaStreamCreate(&streams[i]);
}
// 异步执行多个Kernel
kernel1<<<grid, block, 0, streams[0]>>>();
kernel2<<<grid, block, 0, streams[1]>>>();
kernel3<<<grid, block, 0, streams[2]>>>();
kernel4<<<grid, block, 0, streams[3]>>>();
cudaDeviceSynchronize();