其他分享
首页 > 其他分享> > GPU优化

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名称,……)进行拷贝

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

指令吞吐受限

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

标签:__,warp,访问,内存,GPU,优化,CUDA
来源: https://blog.csdn.net/weixin_51942493/article/details/122869253