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 C 提供了很多实用的 API, CUDA 提供了许多实用的库(eg. cuBlas,cuSparse,Thrust 库)
优化方法:
具体实现:
Texture cache
用__ldg()指定只读缓存
对于只读数据,添加上__ldg()之后,GPU便可以直接从更快的texture缓存中读取。
Shared Memory
使用关键字__shared__即可在GPU中指定一定大小的SM
SM可以让同一个block中的所有线程自由访问,如果多个线程需要访问同一组数据,可以考虑把数据存放在SM中再进行计算。
Constant cache
使用关键字__constant__申明GPU中一个constant cache
使用cudaMemcpyToSymbol(cc名称,……)进行拷贝
加上后缀“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的分析工具,可以得到相关参数)解决。
六 GPU 并行优化的几种典型策略
GPU并行运算与CUDA编程–优化篇
NVIDIA Nsight Systems 入门及使用