GPU架构原理

架构对比

CPU架构（延迟优化）:
┌────────────────────────────────────────────────────┐
│ 控制单元 │ ALU │ 大缓存 │ ALU │ 控制单元 │ ALU │    │
│ (复杂)   │     │        │     │ (复杂)   │     │    │
└────────────────────────────────────────────────────┘
特点：少量强核心，大缓存，复杂控制

GPU架构（吞吐优化）:
┌────────────────────────────────────────────────────┐
│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│
│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│ALU│
└────────────────────────────────────────────────────┘
│                  共享缓存/控制                      │
└────────────────────────────────────────────────────┘
特点：大量弱核心，简单控制，高并行度

SM结构（以H100为例）

一个SM包含：
├── CUDA核心: 128个 (FP32)
├── Tensor Core: 4个 (第4代)
├── SFU: 32个 (特殊功能单元)
├── 共享内存: 228 KB
├── L1缓存: 256 KB
├── 寄存器文件: 256 KB
├── 线程调度器: 4个
└── Warp调度器: 4个

H100共有132个SM，总计16896个CUDA核心

GPU存储层次

速度从快到慢：

寄存器 (Register)
  │ 延迟: 1周期
  │ 容量: 每SM 256KB (H100)
  ▼
共享内存 (Shared Memory)
  │ 延迟: ~20周期
  │ 容量: 每SM 228KB (H100)
  ▼
L1缓存 (L1 Cache)
  │ 延迟: ~30周期
  │ 容量: 每SM 256KB (H100)
  ▼
L2缓存 (L2 Cache)
  │ 延迟: ~200周期
  │ 容量: 50MB (H100)
  ▼
HBM显存 (Global Memory)
  │ 延迟: ~500周期
  │ 容量: 80GB (H100)
  │ 带宽: 3.35 TB/s
  ▼
主机内存 (Host Memory)
  │ 通过PCIe/NVLink访问
  │ 带宽: 64 GB/s (PCIe 5.0) / 900 GB/s (NVLink 4.0)

历代架构对比

Kernel示例

// 定义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];
    }
}

// 主机调用
int main() {
    int n = 1000000;
    float *d_a, *d_b, *d_c;
    
    // 分配显存
    cudaMalloc(&d_a, n * sizeof(float));
    cudaMalloc(&d_b, n * sizeof(float));
    cudaMalloc(&d_c, n * sizeof(float));
    
    // 配置执行参数
    int blockSize = 256;
    int numBlocks = (n + blockSize - 1) / blockSize;
    
    // 启动Kernel
    vectorAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);
    
    // 同步
    cudaDeviceSynchronize();
    
    return 0;
}

CUDA线程层次结构

Grid（网格）
  └── Block（线程块）[最多3维]
        └── Thread（线程）[最多3维]

层次关系：
- Grid: 所有线程的集合，由多个Block组成
- Block: 线程组，Block内线程可同步和共享内存
- Thread: 最小执行单位

限制（H100为例）：
- 每Block最多1024线程
- 每SM最多1536线程
- 每SM最多32个Block
- Grid最大维度: 2^31-1 × 65535 × 65535

CUDA内存类型

// 1. 全局内存 (Global Memory)
__device__ float globalData[1024];  // 所有线程可访问

// 2. 共享内存 (Shared Memory)
__shared__ float sharedData[256];   // Block内线程共享

// 3. 寄存器 (Register)
float regData;  // 线程私有，最快

// 4. 常量内存 (Constant Memory)
__constant__ float constData[256];  // 只读，有缓存

// 5. 纹理内存 (Texture Memory)
texture<float, 1> texRef;  // 特殊读取模式

// 共享内存优化示例
__global__ void matMulShared(float *A, float *B, float *C) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    // 加载到共享内存
    As[ty][tx] = A[row * K + k + tx];
    Bs[ty][tx] = B[(k + ty) * N + col];
    __syncthreads();
    
    // 计算...
}

NVIDIA Nsight工具

# 使用nvprof分析
nvprof ./my_cuda_program

# 使用Nsight Systems
nsys profile ./my_cuda_program

# 使用Nsight Compute分析Kernel
ncu ./my_cuda_program

# 关键指标
- Memory Throughput: 内存吞吐量
- Compute Throughput: 计算吞吐量
- Occupancy: SM占用率
- Warp Execution Efficiency: Warp执行效率
- Memory Access Pattern: 内存访问模式