第3章：统一架构革命 (2006-2009)

传统GPU管线 (2005年前)
┌─────────┐     ┌─────────┐     ┌─────────┐
│ 顶点    │────▶│ 几何    │────▶│ 像素    │
│ 着色器  │     │ 处理器  │     │ 着色器  │
└─────────┘     └─────────┘     └─────────┘
    ▲                ▲                ▲
    │                │                │
固定功能         部分可编程       可编程但受限

问题：
• 硬件利用率低（某些阶段空闲）
• 编程模型复杂（需要映射到图形API）
• 无法进行通用计算

CUDA设计目标
┌────────────────────────────────────────┐
│          易用性 (C/C++扩展)              │
├────────────────────────────────────────┤
│        可扩展性 (硬件抽象层)             │
├────────────────────────────────────────┤
│        高性能 (直接硬件访问)             │
├────────────────────────────────────────┤
│      生态系统 (库、工具、文档)           │
└────────────────────────────────────────┘

硬件现实 → CUDA抽象 → 编程便利

芯片     →  Grid      → 一个内核启动
SM阵列   →  Block     → 协作线程组  
Warp执行 →  Thread    → 独立执行流

关键创新：
• 自动映射：程序员不需要了解硬件细节
• 可扩展性：同样代码在不同GPU上自动扩展
• 同步原语：__syncthreads()实现Block内同步

架构演进对比分析

传统专用架构（GeForce 7900 GTX, 2006）：
┌──────────────────────────────────────┐
│  顶点着色器    像素着色器    几何单元    │
│  ┌──────┐    ┌──────┐    ┌──────┐  │
│  │  x8  │    │ x24  │    │  x4  │  │
│  └──────┘    └──────┘    └──────┘  │
│                                      │
│  利用率统计（实际游戏测试）：           │
│  • 顶点单元：30-50%                  │
│  • 像素单元：60-80%                  │  
│  • 几何单元：10-20%                  │
│  综合效率：40-45%                    │
└──────────────────────────────────────┘

革命性统一架构（G80/Tesla, 2006）：
┌──────────────────────────────────────┐
│        128个统一标量处理器（SP）        │
│  ┌────────────────────────────────┐  │
│  │  动态任务分配调度器              │  │
│  └────────────────────────────────┘  │
│  ↓        ↓        ↓        ↓       │
│  顶点    像素    几何    计算        │
│                                      │
│  利用率提升：                         │
│  • 所有单元：85-95%                  │
│  • 负载自动均衡                      │
│  • 零空闲时间                        │
│  综合效率：90%+                      │
└──────────────────────────────────────┘

实际测试数据（2007年1月）

应用类型        G71(7900GTX)  G80(8800GTX)  提升
─────────────────────────────────────────────
3D游戏平均        100%          185%        1.85×
顶点密集场景       100%          310%        3.1× 
像素密集场景       100%          165%        1.65×
几何处理          100%          420%        4.2×
GPGPU计算         100%          890%        8.9×

G80 流多处理器(SM)结构
┌─────────────────────────────────┐
│      Streaming Multiprocessor    │
├─────────────────────────────────┤
│   ┌─────┐ ┌─────┐ ... ┌─────┐  │
│   │ SP  │ │ SP  │     │ SP  │  │ 8个标量处理器(SP)
│   └─────┘ └─────┘     └─────┘  │
├─────────────────────────────────┤
│        指令调度单元               │
├─────────────────────────────────┤
│   ┌──────────────────────────┐  │
│   │   共享内存 (16KB)         │  │
│   └──────────────────────────┘  │
├─────────────────────────────────┤
│   ┌──────────────────────────┐  │
│   │   寄存器文件 (32KB)       │  │
│   └──────────────────────────┘  │
├─────────────────────────────────┤
│   ┌──────────────────────────┐  │
│   │   常量缓存 (8KB)          │  │
│   └──────────────────────────┘  │
└─────────────────────────────────┘

完整G80：16个SM × 8个SP = 128个CUDA核心

G80完整内存层次结构
┌──────────────────────────────────────────────┐
│  层级        容量      延迟      带宽        作用域  │
├──────────────────────────────────────────────┤
│  寄存器     32KB/SM   1 cyc    8TB/s       线程    │
│  共享内存   16KB/SM   2 cyc    1.3TB/s     Block   │
│  常量缓存   8KB/SM    2 cyc    1.3TB/s     全局    │
│  纹理缓存   8KB/SM    20 cyc   200GB/s     全局    │
│  全局内存   768MB     400 cyc  86.4GB/s    设备    │
└──────────────────────────────────────────────┘

内存访问模式优化
┌──────────────────────────────────────────────┐
│                 合并访问 (Coalescing)               │
│  Warp中32个线程访问连续内存→合并为一次事务         │
│                                                      │
│  好的模式： T0→M[0], T1→M[1], ..., T31→M[31]       │
│  坏的模式： T0→M[0], T1→M[32], ..., T31→M[992]      │
│                                                      │
│  性能差异：32倍                                       │
└──────────────────────────────────────────────┘

CUDA线程层次模型
                Grid
    ┌──────────────────────────┐
    │  ┌────┐ ┌────┐ ┌────┐   │
    │  │B0,0│ │B0,1│ │B0,2│   │
    │  └────┘ └────┘ └────┘   │
    │  ┌────┐ ┌────┐ ┌────┐   │
    │  │B1,0│ │B1,1│ │B1,2│   │
    │  └────┘ └────┘ └────┘   │
    └──────────────────────────┘
              ↓
           Block(1,1)
    ┌──────────────────────────┐
    │ T0 T1 T2 T3 ... T31      │ Warp 0
    │ T32 T33 ... T63          │ Warp 1
    │ ...                      │
    │ T480 ... T511            │ Warp 15
    └──────────────────────────┘

关键概念：
• Warp：32个线程的执行单位（SIMT）
• Block：最多512个线程，可以同步
• Grid：多个Block，无法直接同步

// 完整CUDA程序：向量加法对比

// CPU基准实现
void vectorAdd_CPU(float *a, float *b, float *c, int n) {
    for (int i = 0; i < n; i++) {
        c[i] = a[i] + b[i];
    }
}

// CUDA内核函数
__global__ void vectorAdd_GPU(float *a, float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

// 主程序示例
int main() {
    int n = 1024 * 1024;  // 1M元素
    size_t size = n * sizeof(float);

    // 分配主机内存
    float *h_a = (float*)malloc(size);
    float *h_b = (float*)malloc(size);
    float *h_c = (float*)malloc(size);

    // 分配GPU内存
    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;
    vectorAdd_GPU<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

    // 拷贝结果回主机
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 清理
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);
}

测试环境：

- CPU: Intel Core 2 Duo E6850 @ 3.0GHz
- GPU: NVIDIA GeForce 8800 GTX
- 数据量: 128MB浮点数

性能结果：
┌─────────────────┬───────────┬──────────┐
│ 实现方式          │ 带宽      │ 加速比   │
├─────────────────┼───────────┼──────────┤
│ CPU单线程       │ 0.5 GB/s  │ 1×      │
│ CPU SSE         │ 2.1 GB/s  │ 4.2×    │
│ CUDA未优化      │ 35 GB/s   │ 70×     │
│ CUDA合并访问    │ 80 GB/s   │ 160×    │
└─────────────────┴───────────┴──────────┘

关键优化点：
• 内存合并：连续访问模式
• 高占用率：768线程/SM
• 带宽利用：达到理论值93%

// 使用共享内存优化的矩阵乘法
#define TILE_SIZE 16

__global__ void matrixMul(float *A, float *B, float *C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[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) {
        // 加载到共享内存
        As[ty][tx] = A[row * N + t * TILE_SIZE + tx];
        Bs[ty][tx] = B[(t * TILE_SIZE + ty) * N + col];
        __syncthreads();

        // 计算部分结果
        for (int k = 0; k < TILE_SIZE; ++k)
            sum += As[ty][k] * Bs[k][tx];
        __syncthreads();
    }

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

// 性能：~200 GFLOPS (约40%峰值)

2007-2008年Tesla早期应用领域

石油勘探
├── 地震数据处理：10×加速
├── 储层模拟：15×加速
└── 客户：斯伦贝谢、雪佛龙

金融计算
├── 期权定价：100×加速
├── 风险分析：50×加速
└── 客户：巴克莱、JP摩根

科学研究
├── 分子动力学：20×加速
├── 天体物理：30×加速
└── 客户：斯坦福、橡树岭实验室

医疗影像
├── CT重建：15×加速
├── MRI处理：25×加速
└── 客户：西门子、GE医疗

GeForce 8系列产品定位（2006-2008）

高端发烧级
├── 8800 Ultra (768MB, $829)
├── 8800 GTX (768MB, $599)
└── 8800 GTS (640MB/320MB, $449/$299)

中端主流
├── 8600 GTS (256MB, $229)
├── 8600 GT (256MB, $149)
└── 8500 GT (256MB, $89)

入门级
├── 8400 GS (256MB, $59)
└── 8300 GS (128MB, $39)

移动版本
├── 8800M GTX (笔记本高端)
├── 8600M GT (笔记本主流)
└── 8400M GS (笔记本入门)

架构演进对比
                G80/GT200              Fermi(GF100)
                (2006-2009)            (2010)
┌─────────────────────────┐   ┌─────────────────────────┐
│   图形渲染优化           │   │   科学计算优化           │
│   • 纹理单元为主        │   │   • 双精度浮点          │
│   • 单精度浮点          │   │   • ECC内存              │
│   • 有限缓存            │   │   • 真正缓存层次         │
└─────────────────────────┘   └─────────────────────────┘

关键转变：
• 客户群体：游戏玩家 → HPC用户
• 性能指标：FPS → FLOPS
• 可靠性：偶尔错误可接受 → 零容错

Fermi SM (Streaming Multiprocessor) 架构
┌──────────────────────────────────────────┐
│          Instruction Cache                │
├──────────────────────────────────────────┤
│     Warp Scheduler    Warp Scheduler      │
│     Dispatch Unit     Dispatch Unit       │
├──────────────────────────────────────────┤
│  ┌──────────────┐  ┌──────────────┐     │
│  │  16 CUDA     │  │  16 CUDA     │     │
│  │  Cores       │  │  Cores       │     │ 32个CUDA核心
│  └──────────────┘  └──────────────┘     │
├──────────────────────────────────────────┤
│  ┌──────────────────────────────────┐   │
│  │    16 Load/Store Units            │   │
│  └──────────────────────────────────┘   │
├──────────────────────────────────────────┤
│  ┌──────────────────────────────────┐   │
│  │    4 Special Function Units       │   │ 特殊函数单元
│  └──────────────────────────────────┘   │
├──────────────────────────────────────────┤
│  ┌──────────────────────────────────┐   │
│  │  64KB Shared Memory/L1 Cache      │   │ 可配置缓存
│  │  (48KB/16KB or 16KB/48KB)        │   │
│  └──────────────────────────────────┘   │
├──────────────────────────────────────────┤
│  ┌──────────────────────────────────┐   │
│  │    32K 32-bit Registers           │   │ 寄存器文件
│  └──────────────────────────────────┘   │
└──────────────────────────────────────────┘

完整GF100：16个SM × 32 CUDA核心 = 512核心

浮点性能对比 (GFLOPS)
              单精度(FP32)    双精度(FP64)    比例
G80 (2006)      518            64            8:1
GT200 (2008)    933            78            12:1
Fermi (2010)    1030           515           2:1
              ↑               ↑
        游戏足够          科学计算关键

应用影响：
• 分子动力学：精度要求高
• 气象模拟：累积误差敏感
• 金融建模：小数点精度关键

CUDA版本演进
CUDA 1.0 (2007)          CUDA 3.0 (2010)
├── 基础C扩展            ├── C++类支持
├── 基本内核             ├── 函数指针
├── 纹理内存             ├── 递归
└── 原子操作             ├── printf调试
                        ├── 统一寻址空间
                        └── GPU Direct

开发效率提升：
• 调试时间：减少60%
• 代码复杂度：降低40%
• 移植难度：大幅下降

nForce产品线时间轴
2001 ├── nForce：首款产品，集成GPU
2002 ├── nForce2：AMD平台成功
2004 ├── nForce3：首个AMD64芯片组
2005 ├── nForce4：SLI技术整合
2006 ├── nForce 500：Intel平台扩展
2007 ├── nForce 600：最后辉煌
2008 ├── nForce 700：市场份额下滑
2009 ├── 宣布退出芯片组市场

市场份额变化：
2005年：35%（AMD平台）
2006年：28%
2007年：20%
2008年：15%
2009年：<10%

芯片组团队重新分配（约500人）
├── Tesla/CUDA团队 (200人)
│   └── 加强HPC产品开发
├── Tegra移动团队 (150人)
│   └── ARM SoC开发
├── GPU架构团队 (100人)
│   └── Fermi后续开发
└── 软件工具团队 (50人)
    └── 开发者生态建设

CUDA教学中心（2007-2009）
北美
├── 斯坦福大学：并行计算课程
├── MIT：计算科学应用
├── 伊利诺伊大学：Wen-mei Hwu实验室
├── 哈佛大学：计算物理
└── UC伯克利：计算机架构

欧洲
├── 剑桥大学：科学计算
├── ETH苏黎世：高性能计算
└── INRIA：计算机视觉

亚洲
├── 清华大学：GPU计算中心
├── 东京大学：超算应用
└── IIT印度：并行编程

培养成果：
• 2007年：10所大学，500名学生
• 2008年：50所大学，5000名学生
• 2009年：200所大学，20000名学生

地震数据处理革命
传统CPU集群              CUDA GPU方案
├── 1000节点              ├── 50节点
├── 10MW功耗              ├── 500KW功耗
├── $1000万成本           ├── $100万成本
├── 处理时间：30天        ├── 处理时间：3天
└── 维护复杂              └── 维护简单

关键客户案例：
斯伦贝谢(Schlumberger)
├── 逆时偏移(RTM)：15×加速
├── 全波形反演(FWI)：20×加速
└── 年节省成本：$5000万

雪佛龙(Chevron)
├── 储层模拟：10×加速
├── 地质建模：8×加速
└── 项目周期：6个月→3周

分子动力学模拟 - AMBER
┌────────────────────────────────┐
│  蛋白质折叠模拟（10万原子）      │
├────────────────────────────────┤
│ CPU集群：100ns/天               │
│ 1×Tesla：400ns/天               │
│ 4×Tesla：1500ns/天              │
└────────────────────────────────┘
科研效率：月→天

天体物理 - N体问题
┌────────────────────────────────┐
│  星系碰撞模拟（100万星体）       │
├────────────────────────────────┤
│ CPU：O(N²) = 10¹²次计算         │
│ GPU：树算法+并行 = 30×加速      │
└────────────────────────────────┘

2007年 CUDA 1.0工具
├── nvcc编译器（基础）
├── cuda-gdb（命令行）
└── Visual Profiler（基础）

2009年 CUDA 3.0工具
├── nvcc（C++支持）
├── Parallel Nsight（Visual Studio集成）
├── cuda-memcheck（内存调试）
├── CUDA Profiler（性能分析）
└── GPU占用率计算器

开发效率提升：
• 调试时间：-70%
• 优化周期：-50%
• 学习曲线：显著降低

AMD Stream vs NVIDIA CUDA对比

技术栈对比：
NVIDIA (2007-2009)         AMD (2007-2009)
├── CUDA C/C++             ├── Brook+ (学术语言)
├── 完整工具链             ├── CAL (底层API)
├── 丰富文档               ├── 文档匮乏
├── 大学计划               ├── 有限支持
└── 100+应用案例          └── <10案例

市场结果：
• CUDA开发者：50,000+ (2009)
• Stream开发者：<1,000 (2009)
• 最终AMD在2011年放弃Stream，转向OpenCL

Intel Larrabee (2008-2010，最终取消)
┌─────────────────────────────────┐
│   32个x86核心（基于Pentium）     │
│   512位向量单元                  │
│   一致性缓存                     │
│   x86兼容性                      │
└─────────────────────────────────┘

失败原因：
• 功耗过高：300W+ (vs Tesla 250W)
• 性能不足：图形性能仅GTX 260级别
• 软件栈不成熟：缺乏生态系统
• 定位模糊：既非最佳GPU也非最佳CPU
• 2010年5月宣布取消GPU产品化

市场份额变化 (2006-2009)
        2006    2007    2008    2009
NVIDIA   48%     65%     70%     73%
AMD/ATI  51%     33%     28%     25%
其他      1%      2%      2%      2%

关键产品对比：
2008年高端市场
├── GTX 280：$649，市场领导者
├── HD 4870：$299，性价比选择
└── 结果：NVIDIA利润率高，AMD份额增长有限

2006-2009架构演进
      G80           GT200         Fermi(设计中)
      (2006)        (2008)        (2010)
       │              │              │
    128核心       240核心        512核心
       │              │              │
    无缓存        无缓存         L1+L2缓存
       │              │              │
    单精度      双精度受限      完整双精度
       │              │              │
  CUDA 1.0      CUDA 2.0       CUDA 3.0

CUDA生态系统增长（2007-2009）
┌─────────────────────────────────┐
│ 指标            2007年   2009年   │
├─────────────────────────────────┤
│ 开发者数量       1K      50K      │
│ 应用程序         10      500+     │
│ 大学课程         5       200+     │
│ 科研论文         10      1000+    │
│ CUDA下载量       10K     1M+      │
│ Tesla销售        $1M     $100M+   │
└─────────────────────────────────┘

计算范式转变
串行计算时代              并行计算时代
(1970-2005)              (2006-至今)
    │                        │
单核CPU主导              GPU+CPU异构
    │                        │
摩尔定律驱动            并行度驱动
    │                        │
频率提升                核心数增加
    │                        │
冯诺依曼瓶颈            大规模并行