CUDA C最佳实践-CUDA Best Practices(三)

10. 运行配置优化

10.1. 占用

10.1.1. 计算占用

10.2. 同步Kernel执行

10.3. 多上下文

10.4. 隐藏寄存器依赖

10.5. 线程和线程块启发

10.6. 共享内存的效果

11. 指令优化

知道底层命令是怎么执行的对优化来说很有帮助。不过文档建议要在做过所有高级优化之后再对这进行考虑。

11.1. 算数指令

强烈推荐单精度浮点数!!!

11.1.1. 除/取膜 指令

按位操作永远比普通的操作快,比如当n是2的幂的时候,(i>>log2(n))要比i/n快得多。并且i%n和(i & (n-1))也是相等的。详情查看编程指南

11.1.2. 平方根倒数

要求平方根的倒数使用rsqrtf()

11.1.3. 其他算数指令

要避免double向float的自动转换。我们要在常数后面加f来避免这种事情的发生,因为它会增加多余的时钟周期。并且对于单精度浮点数,建议使用单精度的数学函数和操作。而且在普遍意义上来说,单精度比双精度快。

11.1.4. 小指数取幂

这是啥意思呢,看这个表就知道了:

就是说,在这种情况下,要采用的这种组合的情况而不是直接无脑设置分数。

11.1.5. 数学库

当速度要求超过精度时,使用快速数学库。运行时函数库提供两种类型的函数,__functionName() 和 functionName().后面这种一般比较耗时但是比较精确。而且你还能使用-use_fast_math这种操作让nvcc让后面的转换成前面的,当精度要求不高的时候可以使用这个设置。

另外,当计算类似x^2,x^3这样的整数指数的时候,使用连续相乘会比用pow()函数要开销少。

还有,用 sinpi()替换sin(π*),其他三角函数同理。就是反正有专用的函数要用专用的,别瞎整。

11.1.6. 精度相关的编译标志

nvcc有一些编译开关:

  • ftz=true (非规格化数据转换成零)
  • prec-div=false (精度更低的除法)
  • prec-sqrt=false (精度更低的开平方)
  • -use_fast_math(精度更低的函数)

11.2. 内存指令

尽量避免使用全局内存。尽可能使用共享内存

12. 控制流

12.1. 分支与分歧

一个warp里尽量不要分支。就是一旦遇到分支,warp里的thread要等其他的都运行完才可以。任何控制流指令(if , switch , do , for , while)都能显著影响到指令吞吐量。

12.2. 分支预测

编译器会展开循环或者优化if来进行分支预测。这样的话,warp就不会有分支。程序猿可以使用#pragma unroll来展开循环。想知道更多去看编程指南。

在使用这种分支预测来优化指令时,编译器会给相关于各个线程的指令设置true or false,虽然每个指令都计划被运行,但是实际上只有那些被标记为true的线程执行。这种优化其实是有阈值的,当这个分支的情况少于一定值时会进行替换。

12.3. 有符号或者无符号的循环计数器

用有符号整数做计数器!!!无符号整数溢出被很好地定义,而有符号的没定义,因此编译器就给它优化。比如:

for (i = 0; i < n; i++) {
    out[i] = in[offset + stride*i];
}

有一个stride*i,可能溢出32位整数,因此如果i被定义为无符号数,溢出控制可能会阻止一些诸如截断之类的优化。如果i被定义为有符号整型,编译器就有机会做优化。

12.4. 循环中的线程同步分支

在分支语句中尽量避免使用__syncthreads().

如果在一些分支语句中使用同步函数,可能会造成无法预计的错误(所以到底是什么错误文档也没说)。

unsigned int imax = blockDim.x * ((nelements + blockDim.x - 1)/ blockDim.x);
for (int i = threadidx.x; i < imax; i += blockDim.x)
{
    if (i < nelements)
    {
        ...
    }
    __syncthreads();
    if (i < nelements)
    {
        ...
    }
}

这里的imax被设置成了warp大小的整数倍,可以解决这一问题。所以在使用同步语句的时候一定要注意。可以使用thread_active标志来指出哪些线程是活动的。

13. 实施CUDA应用

优化之后要将实际结果和期望结果比较,再次APOD循环。在进行更深度的优化之前,先把当前的程序部署起来,这样有很多好处,比如允许使用者对当前的应用进行评估,并且减小了应用的风险因为这是一种循序渐进的演化而不是改革。

14. 理解程序运行环境

要注意两点,一是计算能力,二是CUDA运行时和驱动API的版本。

14.1. CUDA计算能力

可以通过CUDA的一个例子deviceQuery来查看计算能力:

14.2. 额外的硬件数据

其他的硬件数据可以通过cudaGetDeviceProperties()这个函数来获得。

14.3. CUDA运行时和驱动API版本

CUDA运行时和驱动API是程序运行的接口。重点是,CUDA的驱动API是后向兼容而不是前向兼容(向后兼容就是新的版本能用旧的接口,旧的版本不能用新的接口):

14.4. 选择哪个运算能力的版本

在编译的时候可以用-arch 来选择计算能力

14.5. CUDA运行时

15. 部署准备

15.1. 测试CUDA可用性

15.2. 错误控制

15.3. 在最大的计算能力下编译

15.4. 分配CUDA运行时和库

15.4.1. CUDA Toolkit Library Redistribution

15.4.1.1. Which Files to Redistribute

15.4.1.2. Where to Install Redistributed CUDA Libraries

16. 实施的基本工具

16.1. Nvidia-SMI

16.1.1. Queryable state

16.1.2. Modifiable state

16.2. NVML

16.3. 集群管理器工具

16.4. 即时编译器缓存管理器工具

16.5. CUDA_VISIBLE_DEVICES

A. 建议和最佳实践

A.1. 优化阶段总结

B. nvcc 编译器参数

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏技术专栏

Spark SQL/Hive调优

任务进度长时间维持在99%(或100%),查看任务监控页面,发现只有少量(1个或几个)reduce子任务未完成。因为其处理的数据量和其他reduce差异过大。 ...

79920
来自专栏YoungGy

MMD_3b_StreamAlgorithms

Stream概述 DBMS的区别 Stream模型 Query种类 应用 Sliding Windows 简介 例子 Bloom Filter motivati...

17790
来自专栏和蔼的张星的图像处理专栏

DSP图像处理

最近着手把CSK移植到DSP中,先看一些DSP中图像处理的一些例子,第一件事当然就是怎么把图像数据倒入CCS工程中了,去年倒是用过一点CCS,再拿起来已经忘得差...

30620
来自专栏吾爱乐享

白盒测试的测试方法及基本路径测试法

22930
来自专栏生信宝典

R语言学习 - 箱线图一步法

箱线图 - 一步绘制 绘图时通常会碰到两个头疼的问题: 有时需要绘制很多的图,唯一的不同就是输入文件,其它都不需要修改。如果用R脚本,需要反复替换文件名,繁琐又...

38650
来自专栏飞雪无情的博客

Go语言实战笔记(二十二)| Go 基准测试

基准测试,是一种测试代码性能的方法,比如你有多种不同的方案,都可以解决问题,那么到底是那种方案性能更好呢?这时候基准测试就派上用场了。

12930
来自专栏生信技能树

转录组数据拼接之应用篇

前前后后接触了一些基因组和转录组拼接的工作,而且后期还会持续进行。期间遇到了各种各样莫名其妙的坑,也尝试了一些不同的方法和软件,简单做一个阶段性小结。上周的今天...

62060
来自专栏轮子工厂

教你用翻译软件快速阅读大量英文文献

对于一些引用的英文文献,我们需要快速地了解整篇文献讲了什么内容,来判断是否可以作为“国内外研究现状”来进行详细分析。

26540
来自专栏SDNLAB

SDN实战团分享(七):YANG模型与OpenDaylight南北向接口

YANG模型是什么? YANG模型是一种数据建模语言,用来建模由NETCONF协议、NETCONF远端过程调用(RPCs)、和NETCONF通知(notific...

66680
来自专栏Python中文社区

Python量子力学计算模拟以及数据可视化

專 欄 ❈Pytlab,Python 中文社区专栏作者。主要从事科学计算与高性能计算领域的应用,主要语言为Python,C,C++。熟悉数值算法(最优化方法,...

1K90

扫码关注云+社区

领取腾讯云代金券