1.3 流处理器（SM/SMX）架构详解

摘要

流处理器（Streaming Multiprocessor，简称SM）是NVIDIA GPU的核心计算单元，是SIMT执行模型的硬件载体。本文深入剖析SM架构的设计原理、内部结构、工作原理及其演进历程，揭示现代GPU如何通过精心设计的SM架构实现大规模并行计算的极致性能。从最初的SM架构到现代的SMX（Streaming Multiprocessor with X），我们将详细分析SM的设计哲学、硬件组成、资源管理机制以及性能优化策略，为理解GPU并行计算提供深入的技术洞察。

1. SM架构基础

1.1 SM的定义与发展历程

1.1.1 SM的概念与作用

SM的定义：
SM是NVIDIA GPU中的基本计算单元，是一个独立的并行处理器，包含多个CUDA核心和丰富的硬件资源。SM负责执行线程块（Thread Block），是GPU并行计算的基本执行单元。

SM的核心作用：

• 并行计算执行：执行CUDA内核，管理多个线程块
• 资源管理：管理计算核心、寄存器、共享内存等资源
• 任务调度：调度和管理线程的执行
• 同步协调：实现线程间的同步和通信

SM在GPU中的位置：

GPU
├── GPU Chip
│   ├── SM 0
│   ├── SM 1
│   ├── SM 2
│   │   ├── CUDA Cores
│   │   ├── Warp Schedulers
│   │   ├── Register File
│   │   ├── Shared Memory
│   │   └── L1 Cache
│   ├── SM 3
│   └── ...
├── Memory Subsystem
│   ├── L2 Cache
│   ├── Memory Controller
│   └── Memory Interfaces
└── Interconnect
    ├── Crossbar Switch
    └── NVLink
​

1.1.2 SM架构的发展历程

第一代SM架构（Tesla架构）：

• 发布时间：2006年（GeForce 8系列）
• 核心特征：
  • 每SM包含8个CUDA核心
  • 单Warp调度器
  • 16KB共享内存
  • 8K寄存器文件
  • 支持CUDA 1.0

第二代SM架构（Fermi架构）：

• 发布时间：2010年（Fermi架构）
• 核心特征：
  • 每SM包含32个CUDA核心
  • 双Warp调度器
  • 64KB共享内存（可配置为48KB共享内存+16KB L1缓存）
  • 64K寄存器文件
  • 支持CUDA 2.0-3.0

第三代SM架构（Kepler架构）：

• 发布时间：2012年（Kepler架构）
• 核心特征：
  • 每SM包含192个CUDA核心
  • 四Warp调度器
  • 64KB-128KB共享内存
  • 64K-128K寄存器文件
  • 支持CUDA 4.0-5.0

第四代SM架构（Maxwell架构）：

• 发布时间：2014年（Maxwell架构）
• 核心特征：
  • 每SM包含128个CUDA核心
  • 四Warp调度器
  • 64KB-96KB共享内存
  • 64K-128K寄存器文件
  • 支持CUDA 6.0-7.0

第五代SM架构（Pascal架构）：

• 发布时间：2016年（Pascal架构）
• 核心特征：
  • 每SM包含128个CUDA核心
  • 四Warp调度器
  • 64KB-96KB共享内存
  • 64K-128K寄存器文件
  • 支持CUDA 8.0-9.0

第六代SM架构（Volta架构）：

• 发布时间：2017年（Volta架构）
• 核心特征：
  • 每SM包含64个CUDA核心 + 8个Tensor Core
  • 四Warp调度器
  • 64KB-96KB共享内存
  • 64K-128K寄存器文件
  • 支持CUDA 9.0-10.0

第七代SM架构（Turing架构）：

• 发布时间：2018年（Turing架构）
• 核心特征：
  • 每SM包含64个CUDA核心 + 8个Tensor Core + 1个RT Core
  • 四Warp调度器
  • 64KB-96KB共享内存
  • 64K-128K寄存器文件
  • 支持CUDA 10.0-11.0

第八代SM架构（Ampere架构）：

• 发布时间：2020年（Ampere架构）
• 核心特征：
  • 每SM包含128个CUDA核心 + 4个Tensor Core + 1个RT Core
  • 四Warp调度器（双发射）
  • 64KB-128KB共享内存
  • 64K-128K寄存器文件
  • 支持CUDA 11.0-12.0

第九代SM架构（Hopper架构）：

• 发布时间：2022年（Hopper架构）
• 核心特征：
  • 每SM包含128个CUDA核心 + 4个Tensor Core + 1个RT Core
  • 四Warp调度器（双发射）
  • 64KB-128KB共享内存
  • 64K-128K寄存器文件
  • 支持FP8精度
  • 支持CUDA 12.0+

第十代SM架构（Blackwell架构）：

• 发布时间：2024年（Blackwell架构）
• 核心特征：
  • 每SM包含128个CUDA核心 + 4个Tensor Core + 1个RT Core
  • 四Warp调度器（四发射）
  • 64KB-128KB共享内存
  • 64K-128K寄存器文件
  • 支持FP8和BF16精度
  • 支持CUDA 12.0+

1.2 SM的基本结构

1.2.1 SM的物理组成

SM的整体结构：

SM (Streaming Multiprocessor)
├── Warp Scheduler Units (2-4)
│   ├── Warp Scheduler 0
│   ├── Warp Scheduler 1
│   ├── Warp Scheduler 2
│   └── Warp Scheduler 3
├── CUDA Cores (64-128)
│   ├── FP32 Unit
│   ├── FP64 Unit
│   ├── INT32 Unit
│   └── Special Function Unit (SFU)
├── Tensor Cores (0-8)
│   ├── Matrix Multiply Unit
│   ├── Activation Unit
│   └── Accumulation Unit
├── RT Core (0-1)
│   ├── Ray Tracing Unit
│   ├── Bounding Volume Hierarchy (BVH) Unit
│   └── Shadow Unit
├── Register File (64K-128K registers)
├── Shared Memory / L1 Cache (64KB-128KB)
├── Constant Cache (64KB)
├── Texture Cache (128KB)
└── Memory Interface
    ├── Load/Store Units
    └── Address Generation Units
​

核心组件详解：

Warp调度器（Warp Scheduler）：

• 数量：2-4个SM
• 功能：调度Warp执行，管理线程执行顺序
• 特点：支持双发射或四发射
• 优化：动态调度，隐藏内存访问延迟

CUDA核心：

• 数量：64-128个SM
• 功能：执行基本的算术逻辑运算
• 类型：FP32、FP64、INT32、SFU
• 特点：简单高效，支持SIMT执行

Tensor Core：

• 数量：0-8个SM
• 功能：执行矩阵乘法运算
• 精度：FP16、TF32、BF16、INT8、INT4
• 特点：高效率矩阵运算，支持混合精度

RT Core：

• 数量：0-1个SM
• 功能：执行光线追踪运算
• 特点：硬件加速光线追踪，支持BVH遍历

寄存器文件（Register File）：

• 容量：64K-128K个32位寄存器
• 功能：存储线程的寄存器数据
• 特点：高速访问，容量有限

共享内存/缓存：

• 容量：64KB-128KB
• 功能：线程块内共享存储和L1缓存
• 特点：高速访问，需要手动管理

1.2.2 SM的资源分配

寄存器分配：

• 每个线程最多分配255个寄存器
• 总寄存器数量：64K-128K（32位）
• 寄存器分配策略：
  • 动态分配：运行时根据需要分配
  • 静态分配：编译时确定寄存器需求
  • 溢出处理：寄存器溢出到本地内存

共享内存分配：

• 总容量：64KB-128KB
• 分配策略：
  • 按Block分配：每个Block的共享内存大小固定
  • 动态调整：共享内存和L1缓存的比例可调
  • 冲突处理：银行冲突检测和解决

执行单元分配：

• CUDA核心：64-128个
• Tensor Core：0-8个
• RT Core：0-1个
• 分配策略：
  • 按指令类型分配
  • 动态负载均衡
  • 优先级调度

1.3 SM的工作原理

1.3.1 SM的启动流程

SM的初始化：

1. GPU启动
   ├── 检测GPU硬件配置
   ├── 初始化SM资源
   ├── 加载固件
   └── 准备执行环境

2. SM激活
   ├── 分配SM资源
   ├── 初始化Warp调度器
   ├── 准备寄存器文件
   └── 加载共享内存

3. 内核启动
   ├── 分配Thread Block
   ├── 分配Warp
   ├── 加载指令缓存
   └── 开始执行
​

Thread Block的分配：

• GPU根据内核的Grid配置分配Thread Block
• 每个Block分配到特定的SM
• Block的Warp数量根据Block大小计算
• 例如：Block大小为256时，分配8个Warp（256÷32=8）

Warp的创建：

• 每个Block被分割为多个Warp
• 每个Warp包含32个线程
• Warp的线程ID为0-31
• Warp的创建和分配对程序员透明

1.3.2 指令执行流程

指令流水线：

取指 → 译码 → 调度 → 发射 → 执行 → 写回
​

详细执行步骤：

1. 取指（Fetch）：

   • 从指令缓存中取出指令
   • 解析指令格式和操作数
   • 确定指令类型和执行单元

2. 译码（Decode）：

   • 解析指令操作码
   • 解析操作数寄存器
   • 确定数据类型和精度

3. 调度（Schedule）：

   • 分配执行单元
   • 检查资源可用性
   • 确定执行优先级

4. 发射（Issue）：

   • 将指令发射到执行单元
   • 设置操作数和数据
   • 启动执行过程

5. 执行（Execute）：

   • 执行实际的计算操作
   • 处理内存访问
   • 执行同步操作

6. 写回（Writeback）：

   • 将结果写回到寄存器
   • 更新程序计数器
   • 处理异常和中断

双发射执行：
在Ampere架构中，SM支持双发射执行：

周期1:
  - 发射指令A到执行单元0
  - 发射指令B到执行单元1
周期2:
  - 发射指令C到执行单元0
  - 发射指令D到执行单元1
​

1.3.3 线程状态管理

线程状态转换：

Active → Waiting → Finished → Exited
  ↓        ↓        ↓        ↓
Ready   Memory    Sync    Complete
Sync   Barrier   Block   Resource
​

线程状态详解：

活跃（Active）：

• 线程正在执行指令
• 使用计算单元和寄存器
• 参与当前指令执行

等待（Waiting）：

• 线程等待内存访问完成
• 等待同步操作完成
• 等待资源可用

完成（Finished）：

• 当前指令执行完成
• 准备执行下一条指令
• 等待调度器调度

退出（Exited）：

• 线程执行完毕
• 释放资源给其他线程
• 不再参与调度

线程掩码管理：

• 每个Warp维护32位掩码
• 掩码控制哪些线程活跃
• 掩码更新指令更新线程状态

2. SMX架构详解

2.1 SMX的设计理念

2.1.1 SMX的定义与特点

SMX的定义：
SMX（Streaming Multiprocessor with X）是NVIDIA在Pascal架构中引入的增强型SM架构，相比之前的SM架构，SMX在性能和效率方面有显著提升。

SMX的核心特点：

• 更高的并行度：支持更多的CUDA核心
• 更高的效率：改进的资源管理机制
• 更强的灵活性：可配置的资源分配
• 更好的编程模型：简化的编程接口

SMX的设计哲学：

• 性能优先：最大化计算吞吐量
• 效率优先：最小化资源浪费
• 可扩展性：支持不同的应用需求
• 兼容性：保持与旧架构的兼容

2.1.2 SMX与SM的对比

硬件对比：

性能对比：

架构对比：

• SM：追求更高的核心数量
• SMX：追求更好的资源利用率和效率
• SMX：支持Tensor Core和RT Core等专用加速器

2.2 SMX的内部结构

2.2.1 SMX的核心组件

SMX的整体结构：

SMX (Streaming Multiprocessor with X)
├── Warp Scheduler Units (4)
│   ├── Dual-Issue Scheduler
│   ├── Instruction Cache
│   └── Resource Management
├── CUDA Cores (128)
│   ├── FP32 Units
│   ├── FP64 Units
│   ├── INT32 Units
│   └── SFU Units
├── Specialized Units
│   ├── Tensor Cores (8, Volta+)
│   ├── RT Cores (1, Turing+)
│   └── Double Precision Units
├── Memory Subsystem
│   ├── Register File (128K registers)
│   ├── Shared Memory / L1 Cache (128KB)
│   ├── Constant Cache (64KB)
│   └── Texture Cache (128KB)
└── Interconnect
    ├── Crossbar Switch
    └── Network Interface
​

Warp调度器增强：

• 双发射调度：每周期可发射2条指令
• 动态优先级：根据资源可用性动态调整优先级
• 指令预取：提前预取指令到缓存
• 分支预测：改进的分支预测算法

CUDA核心增强：

• 更高频率：支持更高的时钟频率
• 混合精度：支持FP16、BF16等混合精度
• 向量化操作：增强的向量化指令支持
• 性能监控：内置性能计数器

专用加速器：

• Tensor Core：硬件加速矩阵运算
• RT Core：硬件加速光线追踪
• 双精度单元：高性能双精度计算

2.2.2 SMX的资源管理

寄存器文件管理：

• 容量扩展：从64K扩展到128K寄存器
• 动态分配：运行时动态分配寄存器
• 溢出处理：寄存器溢出到本地内存
• 优先级调度：基于优先级的寄存器分配

共享内存管理：

• 容量扩展：从64KB扩展到128KB
• 动态配置：共享内存和L1缓存比例可调
• 银行冲突检测：检测和解决内存银行冲突
• 性能优化：优化内存访问模式

执行单元管理：

• 负载均衡：动态分配执行单元
• 优先级调度：基于指令类型的优先级调度
• 资源重用：最大化资源利用率
• 节能模式：根据负载调整功耗

2.3 SMX的工作机制

2.3.1 SMX的调度机制

Warp调度算法：

1. 获取就绪Warp列表
2. 根据优先级选择Warp
3. 检查资源可用性
4. 分配执行单元
5. 发射指令
6. 更新Warp状态
​

优先级策略：

• 内存访问优先级：等待内存的Warp获得高优先级
• 指令缓存优先级：指令在缓存的Warp获得高优先级
• 资源优先级：有足够资源的Warp获得高优先级
• 历史优先级：性能好的Warp获得高优先级

资源检查：

• 寄存器检查：是否有足够寄存器
• 共享内存检查：是否有足够共享内存
• 执行单元检查：是否有对应执行单元
• 内存带宽检查：是否有足够内存带宽

2.3.2 SMX的执行优化

指令流水线优化：

• 流水线深度：优化流水线深度，减少停顿
• 乱序执行：支持指令乱序执行
• 投机执行：支持指令投机执行
• 分支预测：改进的分支预测算法

内存访问优化：

• 合并访问：优化内存访问模式
• 缓存预取：预取数据到缓存
• 带宽优化：最大化内存带宽利用率
• 延迟隐藏：隐藏内存访问延迟

计算优化：

• 向量化：支持向量化计算
• 并行化：最大化并行计算
• 精度优化：支持混合精度计算
• 特殊函数：优化的特殊函数计算

3. SM架构的演进分析

3.1 各代SM架构对比

3.1.1 计算能力演进

第一代到第二代：

• CUDA核心数量：8 → 32（4倍增长）
• 寄存器文件：8K → 64K（8倍增长）
• 共享内存：16KB → 64KB（4倍增长）
• 调度器：单Warp → 双Warp
• 性能提升：约10倍

第二代到第三代：

• CUDA核心数量：32 → 192（6倍增长）
• 寄存器文件：64K → 64K（保持）
• 共享内存：64KB → 128KB（2倍增长）
• 调度器：双Warp → 四Warp
• 性能提升：约3倍

第三代到第四代：

• CUDA核心数量：192 → 128（-33%）
• 寄存器文件：64K → 64K（保持）
• 共享内存：128KB → 64KB（-50%）
• 调度器：四Warp → 四Warp
• 性能提升：约1.2倍

第四代到第五代：

• CUDA核心数量：128 → 128（保持）
• 寄存器文件：64K → 128K（2倍增长）
• 共享内存：64KB → 128KB（2倍增长）
• 调度器：四Warp → 四Warp
• 性能提升：约1.5倍

第五代到第六代：

• 新增Tensor Core：0 → 8个
• 新增双精度单元：支持
• 新增FP16支持：支持
• 性能提升：约2倍（Tensor Core）

第六代到第七代：

• 新增RT Core：0 → 1个
• 新增Tensor Core改进：支持稀疏矩阵
• 新增INT8支持：支持
• 性能提升：约1.5倍

第七代到第八代：

• CUDA核心数量：64 → 128（2倍增长）
• 新增Tensor Core改进：支持TF32
• 新增FP64改进：支持异步执行
• 性能提升：约2倍

第八代到第九代：

• 新增FP8支持：支持
• 新增异步执行：支持
• 新增内存优化：改进
• 性能提升：约1.5倍

3.1.2 内存架构演进

内存带宽演进：

• Tesla：70 GB/s
• Fermi：144 GB/s
• Kepler：320 GB/s
• Maxwell：320 GB/s
• Pascal：480 GB/s
• Volta：900 GB/s
• Turing：616 GB/s
• Ampere：1555 GB/s
• Hopper：3350 GB/s
• Blackwell：5000 GB/s

缓存层次演进：

• L1缓存：16KB → 64KB → 128KB
• L2缓存：768KB → 1.5MB → 6MB → 40MB
• 共享内存：16KB → 64KB → 128KB
• 常量缓存：64KB（保持）
• 纹理缓存：128KB → 128KB → 256KB

3.1.3 编程模型演进

CUDA版本演进：

• CUDA 1.0：基本SIMT支持
• CUDA 2.0：统一内存架构
• CUDA 3.0：动态并行
• CUDA 4.0：GPUDirect
• CUDA 5.0：动态并行改进
• CUDA 6.0：统一内存
• CUDA 7.0：统一内存改进
• CUDA 8.0：Pascal架构支持
• CUDA 9.0：Volta架构支持
• CUDA 10.0：Turing架构支持
• CUDA 11.0：Ampere架构支持
• CUDA 12.0：Hopper架构支持

编程接口演进：

• 基本接口：内核启动、内存管理、同步操作
• 高级接口：流、事件、纹理、统一内存
• 优化接口：共享内存、原子操作、内存合并访问
• 专用接口：Tensor Core、RT Core、双精度运算

3.2 架构演进的关键技术

3.2.1 计算能力的提升

核心数量的优化：

• 早期策略：最大化核心数量
• 中期策略：平衡核心数量和资源
• 后期策略：核心数量趋于稳定，质量提升

执行效率的提升：

• 乱序执行：提高指令执行效率
• 投机执行：减少分支惩罚
• 分支预测：减少分支失败开销
• 指令缓存：减少指令取指延迟

资源利用率的提升：

• 动态调度：提高资源利用率
• 负载均衡：均衡资源使用
• 节能模式：减少闲置功耗
• 性能监控：实时监控资源使用

3.2.2 内存架构的改进

内存带宽的提升：

• 更高的内存频率：提升数据传输速率
• 更多的内存通道：增加并行数据传输
• 更宽的内存位宽：增加单次传输数据量
• 优化的内存控制器：减少内存访问延迟

缓存系统的优化：

• 更大容量的缓存：减少内存访问
• 更智能的缓存管理：提高缓存命中率
• 更细粒度的缓存：提高缓存利用效率
• 优化的缓存替换策略：减少缓存失效

内存访问模式的优化：

• 合并访问：优化内存访问模式
• 缓存预取：预取数据到缓存
• 内存压缩：减少内存占用
• 内存分层：优化数据存储位置

3.2.3 专用加速器的引入

Tensor Core的引入：

• 设计目标：加速矩阵运算
• 硬件实现：专用矩阵乘法单元
• 精度支持：FP16、TF32、BF16、INT8、INT4
• 性能提升：10-20倍矩阵运算性能

RT Core的引入：

• 设计目标：加速光线追踪
• 硬件实现：专用光线追踪单元
• 功能特性：BVH遍历、光线求交、阴影计算
• 性能提升：10-20倍光线追踪性能

双精度单元的增强：

• 设计目标：提高科学计算性能
• 硬件实现：专用双精度浮点单元
• 性能提升：2-3倍双精度运算性能

4. SM架构的性能优化

4.1 硬件层面的优化

4.1.1 SM资源配置优化

寄存器配置优化：

// 优化前：过多的寄存器使用
__global__ void kernel_with_many_registers(float* data, int n) {
    float local_data[100];  // 使用过多寄存器
    for (int i = 0; i < n; i++) {
        local_data[i % 100] = data[i] * 2.0f;
    }
}

// 优化后：减少寄存器使用
__global__ void optimized_kernel(float* data, int n) {
    float temp;  // 使用少量寄存器
    for (int i = 0; i < n; i++) {
        temp = data[i] * 2.0f;
        data[i] = temp;
    }
}
​

共享内存配置优化：

// 优化前：固定共享内存大小
__global__ void fixed_shared_memory_kernel(float* data, int n) {
    __shared__ float shared[64];  // 固定大小
    for (int i = 0; i < 64; i++) {
        shared[i] = data[i];
    }
    __syncthreads();
    for (int i = 0; i < 64; i++) {
        data[i] = shared[i] * 2.0f;
    }
}

// 优化后：动态共享内存配置
__global__ void optimized_shared_memory_kernel(float* data, int n) {
    extern __shared__ float* shared;  // 动态配置
    int tid = threadIdx.x;
    shared[tid] = data[tid];
    __syncthreads();
    data[tid] = shared[tid] * 2.0f;
}
​

4.1.2 执行单元优化

CUDA核心优化：

// 优化前：不充分利用CUDA核心
__global__ void inefficient_kernel(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        data[tid] = sqrtf(data[tid]);  // 单一操作
    }
}

// 优化后：充分利用CUDA核心
__global__ void optimized_kernel(float* data1, float* data2, float* data3, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        // 多个并行操作
        data1[tid] = sqrtf(data1[tid]);
        data2[tid] = powf(data2[tid], 2.0f);
        data3[tid] = expf(data3[tid]);
    }
}
​

Tensor Core优化：

// 优化前：使用基本矩阵乘法
__global__ void naive_matmul(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;
    
    float sum = 0.0f;
    for (int k = 0; k < K; k++) {
        sum += A[row * K + k] * B[k * N + col];
    }
    C[row * N + col] = sum;
}

// 优化后：使用Tensor Core
__global__ void tensor_core_matmul(float* A, float* B, float* C, int M, int N, int K) {
    extern __shared__ float shared_mem[2 * 16 * 16];
    
    float* tile_A = shared_mem;
    float* tile_B = shared_mem + 16 * 16;
    
    int row = blockIdx.y * 16 + threadIdx.y;
    int col = blockIdx.x * 16 + threadIdx.x;
    
    float sum = 0.0f;
    
    for (int k = 0; k < K; k += 16) {
        // 加载tile到共享内存
        for (int i = 0; i < 16; i++) {
            if (row < M && k + i < K) {
                tile_A[threadIdx.y * 16 + i] = A[row * K + k + i];
            }
        }
        
        for (int j = 0; j < 16; j++) {
            if (col < N && k + j < K) {
                tile_B[threadIdx.x * 16 + j] = B[(k + j) * N + col];
            }
        }
        
        __syncthreads();
        
        // 使用Tensor Core进行矩阵乘法
        for (int i = 0; i < 16; i++) {
            for (int j = 0; j < 16; j++) {
                sum += tile_A[threadIdx.y * 16 + i] * tile_B[j * 16 + threadIdx.x];
            }
        }
        
        __syncthreads();
    }
    
    if (row < M && col < N) {
        C[row * N + col] = sum;
    }
}
​

4.1.3 内存访问优化

合并访问优化：

// 优化前：不合并的内存访问
__global__ void uncoalesced_access(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        data[tid * 2] = some_calculation(tid);
        data[tid * 2 + 1] = some_calculation(tid + 1);
    }
}

// 优化后：合并的内存访问
__global__ void coalesced_access(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        data[tid] = some_calculation(tid);
    }
}
​

缓存优化：

// 优化前：直接访问全局内存
__global__ void naive_cache_access(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        float sum = 0.0f;
        for (int i = 0; i < n; i++) {
            sum += data[i];  // 频繁访问全局内存
        }
        data[tid] = sum;
    }
}

// 优化后：使用缓存和共享内存
__global__ void optimized_cache_access(float* data, int n) {
    extern __shared__ float shared_mem[256];
    
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        // 加载数据到共享内存
        for (int i = 0; i < 256; i++) {
            shared_mem[i] = data[tid * 256 + i];
        }
        __syncthreads();
        
        // 使用共享内存计算
        float sum = 0.0f;
        for (int i = 0; i < 256; i++) {
            sum += shared_mem[i];
        }
        
        data[tid] = sum;
    }
}
​

4.2 软件层面的优化

4.2.1 内核优化

块大小优化：

// 优化前：不合适的块大小
dim3 block(16, 16);  // 256个线程
dim3 grid((n + 255) / 256, (m + 255) / 256);

// 优化后：合适的块大小
dim3 block(32, 8);  // 256个线程，更好的内存访问模式
dim3 grid((n + 255) / 256, (m + 255) / 256);
​

线程束优化：

// 优化前：避免Warp分歧
__global__ void divergent_kernel(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        if (data[tid] > 0.0f) {
            data[tid] = sqrtf(data[tid]);  // 分歧
        } else {
            data[tid] = data[tid] * 2.0f;  // 分歧
        }
    }
}

// 优化后：避免Warp分歧
__global__ void optimized_kernel(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        unsigned int mask = __ballot_sync(0xFFFFFFFF, data[tid] > 0.0f);
        if (mask) {
            // 处理正数
            int idx = __ffs_sync(mask) - 1;
            data[tid] = sqrtf(data[tid]);
        }
        if (~mask) {
            // 处理非正数
            int idx = __ffs_sync(~mask) - 1;
            data[tid] = data[tid] * 2.0f;
        }
    }
}
​

同步优化：

// 优化前：频繁同步
__global__ void
​