GPU优化

Total

GPU 并行编程技术，对现有的程序进行并行优化

先对数据集进行分解，然后将任务进行分解

1. 从矩阵角度(数据集)来分析数据，
2. 将输入集和输出集中各个格点的对应关系找出来,
3. 后分派给各个块，各个线程。

识别代码的热点(热点分析)

使用分析工具来找出瓶颈(eg. CUDA Profiler or Parallel Nsight)

使用Nsight Systems分析GPU性能
NVIDIA Nsight Systems 简称nsys，低开销的系统分析工具 存在三种不同的活动区：

1. 分析—收集任何性能数据 （采样和跟踪）
2. 采样—定期停止配置文件（统计了解函数费时）
3. 跟踪-收集有关概要文件或系统中发生的各种活动，(函数调用确切时间和持续时间)

ps. GPU 进行数据处理的时候，CPU 可以考虑其他的运行（将 CPU 并行和 GPU 并行结合）

合理利用内存

特别要注意共享内存的使用（速度接近一级缓存）
必要时对多个内核函数进行融合（注意其中隐含的同步问题）
对于数据的访问采取合并访问的方式 - 使用 cudaMalloc 函数。（充分地利用显卡的带宽，一次访问的数据应当大于 128 字节）

传输过程的优化

内存和显存之间进行数据交换耗时
采用锁页内存（该区域内存和显卡的传递不需要 CPU 来干预）的方式使用主机端内存
锁页内存即调用 cudaHostAlloc 函数。
ps.零复制内存

科学的计算网格

设定维数，块数，块内线程数来实现合并的内存访问，保证最大的内存带宽
多维度的计算网格

CUDA 并行编程

为核心, CUDA C 提供了很多实用的 API, CUDA 提供了许多实用的库(eg. cuBlas,cuSparse,Thrust 库)

Problems

内存带宽受限

优化方法：

1. 用其他内存分担压力 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名称,……)进行拷贝

2. 改变访问顺序，降低上一级内存的cache miss，缓解当前内存的压力
3. 用算法压缩数据/改变数据访问方式(降低不必要的数据访问)

指令吞吐受限

加上后缀“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 入门及使用