GPU优化
作者:互联网
Total
GPU 并行编程技术,对现有的程序进行并行优化
先对数据集进行分解,然后将任务进行分解
- 从矩阵角度(数据集)来分析数据,
- 将输入集和输出集中各个格点的对应关系找出来,
- 后分派给各个块,各个线程。
识别代码的热点(热点分析)
使用分析工具来找出瓶颈(eg. CUDA Profiler or Parallel Nsight)
使用Nsight Systems分析GPU性能
NVIDIA Nsight Systems 简称nsys,低开销的系统分析工具 存在三种不同的活动区:
- 分析—收集任何性能数据 (采样和跟踪)
- 采样—定期停止配置文件(统计了解函数费时)
- 跟踪-收集有关概要文件或系统中发生的各种活动,(函数调用确切时间和持续时间)
ps. GPU 进行数据处理的时候,CPU 可以考虑其他的运行(将 CPU 并行和 GPU 并行结合)
合理利用内存
特别要注意共享内存的使用(速度接近一级缓存)
必要时对多个内核函数进行融合(注意其中隐含的同步问题)
对于数据的访问采取合并访问的方式 - 使用 cudaMalloc 函数。(充分地利用显卡的带宽,一次访问的数据应当大于 128 字节)
传输过程的优化
内存和显存之间进行数据交换耗时
采用锁页内存(该区域内存和显卡的传递不需要 CPU 来干预)的方式使用主机端内存
锁页内存即调用 cudaHostAlloc 函数。
ps.零复制内存
科学的计算网格
设定维数,块数,块内线程数来实现合并的内存访问,保证最大的内存带宽
多维度的计算网格
CUDA 并行编程
为核心, CUDA C 提供了很多实用的 API, CUDA 提供了许多实用的库(eg. cuBlas,cuSparse,Thrust 库)
Problems
内存带宽受限
优化方法:
- 用其他内存分担压力 eg. TEX/Shared Memory/Constant Memory
具体实现:
Texture cache
用__ldg()指定只读缓存
对于只读数据,添加上__ldg()之后,GPU便可以直接从更快的texture缓存中读取。
Shared Memory
使用关键字__shared__即可在GPU中指定一定大小的SM
SM可以让同一个block中的所有线程自由访问,如果多个线程需要访问同一组数据,可以考虑把数据存放在SM中再进行计算。
Constant cache
使用关键字__constant__申明GPU中一个constant cache
使用cudaMemcpyToSymbol(cc名称,……)进行拷贝
- 改变访问顺序,降低上一级内存的cache miss,缓解当前内存的压力
- 用算法压缩数据/改变数据访问方式(降低不必要的数据访问)
指令吞吐受限
加上后缀“f”:
默认情况下,对于立即数“1.0”是双精度的浮点数,加上后缀“f”后,“1.0f”会被编译为单精度浮点数,计算快速。
intrinsic function:
fun()精度高,速度慢;__fun()精度低,速度快。
减少Bank conflict
减少warp里指令的发散
延迟受限
增加active warp的数量, 增加warp来隐藏指令或者data I/O所带来的延迟.
occupancy(用于衡量active warp) = active warp数/SM中所允许的最大warp数
延迟源头(指令的读取/执行的相关性/同步/数据请求/texture等等,利用CUDA的分析工具,可以得到相关参数)解决。
资料
标签:__,warp,访问,内存,GPU,优化,CUDA 来源: https://blog.csdn.net/weixin_51942493/article/details/122869253