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
本站声明： 1. iCode9 技术分享网（下文简称本站）提供的所有内容，仅供技术学习、探讨和分享； 2. 关于本站的所有留言、评论、转载及引用，纯属内容发起人的个人观点，与本站观点和立场无关； 3. 关于本站的所有言论和文字，纯属内容发起人的个人观点，与本站观点和立场无关； 4. 本站文章均是网友提供，不完全保证技术分享内容的完整性、准确性、时效性、风险性和版权归属；如您发现该文章侵犯了您的权益，可联系我们第一时间进行删除； 5. 本站为非盈利性的个人网站，所有内容不会用来进行牟利，也不会利用任何形式的广告来间接获益，纯粹是为了广大技术爱好者提供技术内容和技术思想的分享性交流网站。