- Published on
CUDA 内存类型及特性总结
1. 寄存器(Registers) • 作用:存储线程私有变量(自动变量)。
• 特点:
• 速度最快,延迟最低(1个时钟周期)。
• 每个线程独立拥有自己的寄存器。
• 数量有限(每个SM约64K个32-bit寄存器)。
• 适用场景:频繁访问的局部变量。
2. 本地内存(Local Memory) • 作用:存储线程私有数据(如数组、大结构体)。
• 特点:
• 实际位于全局内存,速度较慢。
• 编译器自动将无法放入寄存器的变量放入本地内存。
• 适用场景:较大的临时变量或动态索引数组。
3. 共享内存(Shared Memory) • 作用:块内线程共享数据(__shared__
)。
• 特点:
• 速度接近寄存器(延迟约10个时钟周期)。
• 每个SM约64KB(可配置为L1缓存)。
• Bank Conflict:32个Bank,每个Bank 4字节宽。
• 适用场景:线程协作、数据复用(如矩阵乘法)。
4. 全局内存(Global Memory) • 作用:所有线程可访问的主存(cudaMalloc
分配)。
• 特点:
• 延迟高(400+时钟周期),带宽高。
• 需要合并访问(Coalesced Access)优化性能。
• 支持__ldg
只读缓存。
• 适用场景:输入/输出大数据存储。
5. 常量内存(Constant Memory) • 作用:存储只读常量(__constant__
)。
• 特点:
• 缓存加速(约8KB缓存)。
• 适合所有线程读取相同值(如滤波器系数)。
• 适用场景:广播式数据(如数学常数)。
6. 纹理内存(Texture Memory) • 作用:优化多维数据的只读访问。
• 特点:
• 自动缓存,支持插值、边界处理。
• 适合空间局部性访问(如图像处理)。
• 适用场景:图像处理、插值计算。
7. 固定内存(Pinned Memory) • 作用:主机端锁定内存(cudaMallocHost
)。
• 特点:
• 加速主机-设备数据传输(避免复制)。
• 不可交换(避免被OS换出)。
• 适用场景:高频主机-设备数据传输。
8. 统一内存(Unified Memory) • 作用:简化主机-设备内存管理(cudaMallocManaged
)。
• 特点:
• 自动迁移数据,编程简单。
• 性能可能低于手动管理。
• 适用场景:简化开发,适合复杂数据结构。
内存性能对比
内存类型 | 延迟 | 带宽 | 作用域 | 优化要点 |
---|---|---|---|---|
寄存器 | 1 cycle | ∞ | 线程私有 | 尽量多用 |
共享内存 | ~10 cycles | 高 | 线程块共享 | 避免Bank Conflict |
常量内存 | ~20 cycles | 中 | 全局只读 | 适合广播数据 |
纹理内存 | ~20 cycles | 中 | 全局只读 | 适合空间局部性 |
全局内存 | 400+ cycles | 高 | 全局可读写 | 合并访问,对齐 |
本地内存 | 同全局内存 | 低 | 线程私有 | 尽量避免使用 |
关键优化策略
最大化寄存器使用:减少本地内存访问。
共享内存优化: • Padding避免Bank Conflict(如
[32][33]
)。• 用于数据复用(如矩阵分块)。
全局内存合并访问: • 连续线程访问连续地址(步长=1最优)。
常量/纹理内存:利用缓存加速只读数据。
统一内存:简化代码,但需测试性能。
一句话总结 • 寄存器 > 共享内存 > 常量/纹理 > 全局内存
• 优化核心:减少全局内存访问,用共享内存/寄存器替代,避免Bank Conflict和未合并访问。
THE END